Public/koboldcpp
1
0
Fork 0
mirror of https://github.com/LostRuins/koboldcpp.git synced 2026-09-18 16:55:14 +02:00
Code Issues Packages Projects Releases Wiki Activity

Merge branch 'upstream' into concedo_experimental

Browse Source 
# Conflicts:
#	.github/actions/ccache-clear/action.yml
#	.github/workflows/build-apple.yml
#	.github/workflows/build-cpu.yml
#	.github/workflows/build-cuda-ubuntu.yml
#	.github/workflows/build-opencl.yml
#	.github/workflows/build-openvino.yml
#	.github/workflows/build-sycl.yml
#	.github/workflows/build-vulkan.yml
#	.github/workflows/build-wasm.yml
#	.github/workflows/build-webgpu.yml
#	.github/workflows/hip-quality-check.yml
#	.github/workflows/server.yml
#	CONTRIBUTING.md
#	README.md
#	ci/run.sh
#	common/CMakeLists.txt
#	common/chat.cpp
#	docs/autoparser.md
#	ggml/src/ggml-webgpu/wgsl-shaders/flash_attn.wgsl
#	ggml/src/ggml-webgpu/wgsl-shaders/flash_attn_tile.wgsl
#	ggml/src/ggml-webgpu/wgsl-shaders/flash_attn_vec_split.wgsl
#	scripts/sync_vendor.py
#	tests/CMakeLists.txt
#	tests/peg-parser/test-json-serialization.cpp
#	tests/peg-parser/tests.h
#	tests/test-chat-auto-parser.cpp
#	tests/test-chat-peg-parser.cpp
#	tests/test-chat-template.cpp
#	tests/test-chat.cpp
#	tests/test-grammar-integration.cpp
#	tests/test-jinja.cpp
#	tests/test-json-schema-to-grammar.cpp
#	tests/test-llama-archs.cpp
#	tests/test-model-resolution.cpp
#	tests/test-recurrent-state-rollback.cpp
#	tools/CMakeLists.txt
This commit is contained in:
Concedo
2026-08-25 20:44:34 +08:00
parent 3265983620 7584430716
commit 4447017602
124 changed files with 3564 additions and 1659 deletions
Download Patch File Download Diff File Expand all files Collapse all files
CMakeLists.txt
+2
View File
@@ -469,6 +469,8 @@ set_target_properties(ggml_v3 PROPERTIES POSITION_INDEPENDENT_CODE ON)
add_library(common2
common/common.cpp
common/common.h
common/json.cpp
common/json.h
common/speculative.cpp
common/speculative.h
common/sampling.cpp
Makefile
+6 -4
View File
@@ -110,10 +110,10 @@ endif
CUBLASLD_FLAGS =
CUBLAS_OBJS =
OBJS_FULL += ggml-alloc.o ggml-cpu-traits.o ggml-quants.o ggml-cpu-quants.o kcpp-quantmapper.o kcpp-repackmapper.o unicode.o unicode-common.o unicode-data.o ggml-threading.o ggml-cpu-cpp.o gguf.o sgemm.o common.o speculative.o llama-impl.o sampling.o budget.o kcpputils.o kcppllmutils.o mtmdaudio.o ggml-rpc.o transport.o hash.o
OBJS_SIMPLE += ggml-alloc.o ggml-cpu-traits.o ggml-quants_noavx2.o ggml-cpu-quants.o kcpp-quantmapper_noavx2.o kcpp-repackmapper_noavx2.o unicode.o unicode-common.o unicode-data.o ggml-threading.o ggml-cpu-cpp.o gguf.o sgemm_noavx2.o common.o speculative.o llama-impl.o sampling.o budget.o kcpputils.o kcppllmutils.o mtmdaudio.o ggml-rpc.o transport.o hash.o
OBJS_SIMPLER += ggml-alloc.o ggml-cpu-traits.o ggml-quants_noavx1.o ggml-cpu-quants.o kcpp-quantmapper_noavx1.o kcpp-repackmapper_noavx1.o unicode.o unicode-common.o unicode-data.o ggml-threading.o ggml-cpu-cpp.o gguf.o sgemm_noavx1.o common.o speculative.o llama-impl.o sampling.o budget.o kcpputils.o kcppllmutils.o mtmdaudio.o ggml-rpc.o transport.o hash.o
OBJS_FAILSAFE += ggml-alloc.o ggml-cpu-traits.o ggml-quants_failsafe.o ggml-cpu-quants.o kcpp-quantmapper_failsafe.o kcpp-repackmapper_failsafe.o unicode.o unicode-common.o unicode-data.o ggml-threading.o ggml-cpu-cpp.o gguf.o sgemm_failsafe.o common.o speculative.o llama-impl.o sampling.o budget.o kcpputils.o kcppllmutils.o mtmdaudio.o ggml-rpc.o transport.o hash.o
OBJS_FULL += ggml-alloc.o ggml-cpu-traits.o ggml-quants.o ggml-cpu-quants.o kcpp-quantmapper.o kcpp-repackmapper.o unicode.o unicode-common.o unicode-data.o ggml-threading.o ggml-cpu-cpp.o gguf.o sgemm.o common.o common-json.o speculative.o llama-impl.o sampling.o budget.o kcpputils.o kcppllmutils.o mtmdaudio.o ggml-rpc.o transport.o hash.o
OBJS_SIMPLE += ggml-alloc.o ggml-cpu-traits.o ggml-quants_noavx2.o ggml-cpu-quants.o kcpp-quantmapper_noavx2.o kcpp-repackmapper_noavx2.o unicode.o unicode-common.o unicode-data.o ggml-threading.o ggml-cpu-cpp.o gguf.o sgemm_noavx2.o common.o common-json.o speculative.o llama-impl.o sampling.o budget.o kcpputils.o kcppllmutils.o mtmdaudio.o ggml-rpc.o transport.o hash.o
OBJS_SIMPLER += ggml-alloc.o ggml-cpu-traits.o ggml-quants_noavx1.o ggml-cpu-quants.o kcpp-quantmapper_noavx1.o kcpp-repackmapper_noavx1.o unicode.o unicode-common.o unicode-data.o ggml-threading.o ggml-cpu-cpp.o gguf.o sgemm_noavx1.o common.o common-json.o speculative.o llama-impl.o sampling.o budget.o kcpputils.o kcppllmutils.o mtmdaudio.o ggml-rpc.o transport.o hash.o
OBJS_FAILSAFE += ggml-alloc.o ggml-cpu-traits.o ggml-quants_failsafe.o ggml-cpu-quants.o kcpp-quantmapper_failsafe.o kcpp-repackmapper_failsafe.o unicode.o unicode-common.o unicode-data.o ggml-threading.o ggml-cpu-cpp.o gguf.o sgemm_failsafe.o common.o common-json.o speculative.o llama-impl.o sampling.o budget.o kcpputils.o kcppllmutils.o mtmdaudio.o ggml-rpc.o transport.o hash.o
# OS specific
ifeq ($(UNAME_S),Linux)
@@ -673,6 +673,8 @@ llama-model.o: src/llama-model.cpp src/llama-model.h src/models/models.h ggml/in
$(CXX) $(CXXFLAGS) -c $< -o $@
common.o: common/common.cpp common/common.h common/log.h
$(CXX) $(CXXFLAGS) -c $< -o $@
common-json.o: common/json.cpp common/json.h
$(CXX) $(CXXFLAGS) -c $< -o $@
speculative.o: common/speculative.cpp common/speculative.h
$(CXX) $(CXXFLAGS) -c $< -o $@
sampling.o: common/sampling.cpp common/common.h common/sampling.h common/log.h
common/arg.cpp
+3 -4
View File
@@ -5,6 +5,7 @@
#include "common.h"
#include "download.h"
#include "json-schema-to-grammar.h"
#include "json.h"
#include "llama.h"
#include "log.h"
#include "sampling.h"
@@ -22,9 +23,6 @@
#include <shellapi.h>
#endif
#define JSON_ASSERT GGML_ASSERT
#include <nlohmann/json.hpp>
#include <algorithm>
#include <cinttypes>
#include <climits>
@@ -33,6 +31,7 @@
#include <filesystem>
#include <fstream>
#include <list>
#include <numeric>
#include <regex>
#include <set>
#include <string>
@@ -56,7 +55,7 @@
#define LLAMA_MAX_URL_LENGTH 2084 // Maximum URL Length in Chrome: 2083
using json = nlohmann::ordered_json;
using json = common_json;
using namespace common_arg_utils;
static std::initializer_list<enum llama_example> mmproj_examples = {
common/chat-auto-parser-generator.cpp
+2 -3
View File
@@ -5,13 +5,12 @@
#include "common.h"
#include "json-schema-to-grammar.h"
#include "log.h"
#include "nlohmann/json.hpp"
#include "peg-parser.h"
#include <stdexcept>
#include <string>
using json = nlohmann::ordered_json;
using json = common_json;
// Helper to iterate over tools/functions
static void foreach_function(const json & tools, const std::function<void(const json &)> & fn) {
@@ -391,7 +390,7 @@ common_peg_parser analyze_tools::build_tool_parser_tag_tagged(parser_build_conte
std::set<std::string> required;
if (params.contains("required")) {
params.at("required").get_to(required);
required = params.at("required").get<std::set<std::string>>();
}
auto schema_info = common_schema_info();
common/chat-auto-parser-helpers.cpp
-3
View File
@@ -4,14 +4,11 @@
#include "chat-peg-parser.h"
#include "chat.h"
#include "log.h"
#include "nlohmann/json.hpp"
#include "peg-parser.h"
#include <cctype>
#include <numeric>
using json = nlohmann::ordered_json;
std::string trim_whitespace(const std::string & str) {
size_t start = 0;
while (start < str.length() && std::isspace(static_cast<unsigned char>(str[start]))) {
common/chat-auto-parser.h
+2 -2
View File
@@ -4,7 +4,7 @@
#include "common.h"
#include "jinja/caps.h"
#include "peg-parser.h"
#include "nlohmann/json.hpp"
#include "json.h"
#include <chrono>
#include <optional>
@@ -12,7 +12,7 @@
#include <utility>
#include <vector>
using json = nlohmann::ordered_json;
using json = common_json;
class common_chat_peg_builder;
common/chat-diff-analyzer.cpp
+3 -3
View File
@@ -4,11 +4,11 @@
#include "chat.h"
#include "common.h"
#include "log.h"
#include "nlohmann/json.hpp"
#include "peg-parser.h"
#include <algorithm>
#include <cctype>
#include <numeric>
#include <ostream>
#include <sstream>
@@ -17,7 +17,7 @@
#define ANSI_ORANGE "\033[1m\x1b[38;5;214m"
#define ANSI_RED "\033[1m\x1b[38;5;196m"
using json = nlohmann::ordered_json;
using json = common_json;
namespace autoparser {
@@ -929,7 +929,7 @@ void analyze_tools::analyze_tool_call_format_json_native(const std::string & cle
int json_end = clean_haystack.find_last_of('}');
std::string cut = clean_haystack.substr(json_start, json_end - json_start + 1);
json call_struct = json::parse(cut);
auto register_field = [&](const std::string & prefix, const nlohmann::detail::iteration_proxy_value<json::iterator> & subel) {
auto register_field = [&](const std::string & prefix, const common_json_entry & subel) {
if (subel.value().is_string() && std::string(subel.value()).find("call0000") != std::string::npos) {
format.id_field = !prefix.empty() ? prefix + "." + subel.key() : subel.key();
} else if (subel.value().is_string() && std::string(subel.value()) == fun_name_needle) {
common/chat-peg-parser.cpp
+1 -3
View File
@@ -4,12 +4,10 @@
#include "ggml.h"
#include "peg-parser.h"
#include <nlohmann/json.hpp>
#include <cstdint>
#include <functional>
using ordered_json = nlohmann::ordered_json;
using ordered_json = common_json;
static std::string_view trim_trailing_space(std::string_view sv, int max = -1) {
int count = 0;
common/chat-peg-parser.h
+6 -6
View File
@@ -128,7 +128,7 @@ class common_chat_peg_builder : public common_peg_parser_builder {
// parameters_order: order in which JSON fields should be parsed
common_peg_parser standard_json_tools(const std::string & section_start,
const std::string & section_end,
const nlohmann::ordered_json & tools,
const common_json & tools,
bool parallel_tool_calls,
bool force_tool_calls,
const std::string & name_key = "",
@@ -143,13 +143,13 @@ class common_chat_peg_builder : public common_peg_parser_builder {
// Legacy-compatible helper for building XML/tagged style tool calls
// Used by tests and manual parsers
common_peg_parser standard_constructed_tools(const std::map<std::string, std::string> & markers,
const nlohmann::ordered_json & tools,
const common_json & tools,
bool parallel_tool_calls,
bool force_tool_calls);
// Helper for Python-style function call format: name(arg1="value1", arg2=123)
// Used by LFM2 and similar templates
common_peg_parser python_style_tool_calls(const nlohmann::ordered_json & tools,
common_peg_parser python_style_tool_calls(const common_json & tools,
bool parallel_tool_calls,
bool allow_json_literals);
@@ -158,19 +158,19 @@ class common_chat_peg_builder : public common_peg_parser_builder {
common_peg_parser python_or_json_value();
// Implementation helpers for standard_json_tools — one per JSON tool call layout mode
common_peg_parser build_json_tools_function_is_key(const nlohmann::ordered_json & tools,
common_peg_parser build_json_tools_function_is_key(const common_json & tools,
const std::string & args_key,
const std::string & effective_args_key,
const std::string & call_id_key,
const std::string & gen_call_id_key);
common_peg_parser build_json_tools_nested_keys(const nlohmann::ordered_json & tools,
common_peg_parser build_json_tools_nested_keys(const common_json & tools,
const std::string & effective_name_key,
const std::string & effective_args_key,
const std::string & call_id_key,
const std::string & gen_call_id_key);
common_peg_parser build_json_tools_flat_keys(const nlohmann::ordered_json & tools,
common_peg_parser build_json_tools_flat_keys(const common_json & tools,
const std::string & effective_name_key,
const std::string & effective_args_key,
const std::string & call_id_key,
common/chat.cpp
+19 -17
View File
@@ -6,6 +6,7 @@
#include "common.h"
#include "ggml.h"
#include "json-schema-to-grammar.h"
#include "json.h"
#include "log.h"
#include "reasoning-budget.h"
#include "chat-auto-parser-generator.cpp"
@@ -30,6 +31,7 @@
#include <ctime>
#include <exception>
#include <functional>
#include <iomanip>
#include <map>
#include <optional>
@@ -39,7 +41,7 @@
#include <utility>
#include <vector>
using json = nlohmann::ordered_json;
using json = common_json;
static std::string format_time(const std::chrono::system_clock::time_point & now, const std::string & format) {
auto time = std::chrono::system_clock::to_time_t(now);
@@ -57,7 +59,7 @@ static json safe_args_parse(const std::string & to_parse) {
}
try {
return json::parse(stripped);
} catch (json::exception & e) {
} catch (const common_json_error & e) {
return stripped;
}
}
@@ -497,17 +499,17 @@ struct messages_inp_normalizer {
json normalized = json::array();
for (const auto & msg : messages) {
json copy = msg;
auto it = copy.find("content");
if (it != copy.end()) {
if (only_typed && it->is_string()) {
*it = json::array({
if (copy.contains("content")) {
json & it = copy.at("content");
if (only_typed && it.is_string()) {
it = json::array({
json{
{"type", "text"},
{"text", it->get<std::string>()},
{"text", it.get<std::string>()},
}
});
} else if (only_string && it->is_array()) {
*it = concat_content_parts(*it);
} else if (only_string && it.is_array()) {
it = concat_content_parts(it);
}
}
normalized.push_back(std::move(copy));
@@ -622,7 +624,7 @@ std::vector<common_chat_tool> common_chat_tools_parse_oaicompat(const json & too
#include "peg-parser.cpp"
#include "chat-peg-parser.cpp"
common_chat_continuation common_chat_continuation_parse(const nlohmann::ordered_json & value) {
common_chat_continuation common_chat_continuation_parse(const common_json & value) {
if (value.is_boolean() && value.get<bool>()) {
return COMMON_CHAT_CONTINUATION_AUTO;
}
@@ -934,7 +936,7 @@ static void foreach_parameter(const json &
const auto & props = params.at("properties");
std::set<std::string> required;
if (params.contains("required") && params.at("required").is_array()) {
params.at("required").get_to(required);
required = params.at("required").get<std::set<std::string>>();
}
for (const auto & [name, prop] : props.items()) {
bool is_required = (required.find(name) != required.end());
@@ -951,7 +953,7 @@ static std::string common_chat_template_direct_apply_impl(
jinja::context ctx(tmpl.source());
// messages_override is already built for this template, do not touch its content parts
nlohmann::ordered_json inp = nlohmann::ordered_json{
json inp = json{
{"messages", messages_override.has_value()
? *messages_override
: messages_inp_normalizer(tmpl.original_caps()).normalize(inputs.messages)},
@@ -1072,7 +1074,7 @@ static common_chat_params common_chat_params_init_ministral_3(const common_chat_
});
} else if (msg.at("content").is_array()) {
auto blocks = msg.at("content");
content.insert(content.end(), blocks.begin(), blocks.end());
content.insert(blocks);
}
}
@@ -2252,7 +2254,7 @@ static common_chat_params common_chat_params_init_deepseek_v3_2(const common_cha
std::set<std::string> required;
if (params.contains("required")) {
params.at("required").get_to(required);
required = params.at("required").get<std::set<std::string>>();
}
auto schema_info = common_schema_info();
@@ -2874,7 +2876,7 @@ static common_chat_params common_chat_params_init_minimax_m3(const common_chat_t
std::set<std::string> required;
if (schema.contains("required")) {
schema.at("required").get_to(required);
required = schema.at("required").get<std::set<std::string>>();
}
std::vector<common_peg_parser> required_elements;
@@ -2986,10 +2988,10 @@ static void system_message_not_supported(json & messages) {
auto & second_msg = messages[1];
second_msg["content"] = first_msg.at("content").get<std::string>()
+ "\n" + second_msg.at("content").get<std::string>();
messages.erase(messages.begin());
messages.erase(0);
} else {
LOG_WRN("Removing system prompt due to template not supporting system role\n");
messages.erase(messages.begin());
messages.erase(0);
}
}
}
common/chat.h
+9 -10
View File
@@ -8,7 +8,7 @@
#include "jinja/runtime.h"
#include "jinja/caps.h"
#include "nlohmann/json_fwd.hpp"
#include "json.h"
#include <chrono>
#include <functional>
@@ -17,7 +17,6 @@
#include <vector>
using chat_template_caps = jinja::caps;
using json = nlohmann::ordered_json;
struct common_chat_templates;
@@ -87,7 +86,7 @@ struct common_chat_msg {
std::string tool_name;
std::string tool_call_id;
nlohmann::ordered_json to_json_oaicompat(bool concat_typed_text = false) const;
common_json to_json_oaicompat(bool concat_typed_text = false) const;
std::string render_content(const std::string & delimiter = "\n\n") const;
@@ -211,7 +210,7 @@ struct common_chat_msg_delimiters {
// split tokens into message spans. skips maps a start index to a length of a region to jump over without matching
common_chat_msg_spans split(const llama_tokens & tokens, const std::map<size_t, size_t> & skips = {}) const;
nlohmann::ordered_json to_json() const;
common_json to_json() const;
};
struct common_chat_tool {
@@ -350,16 +349,16 @@ common_chat_tool_choice common_chat_tool_choice_parse_oaicompat(const std::strin
bool common_chat_templates_support_enable_thinking(const common_chat_templates * chat_templates);
// Parses a JSON array of messages in OpenAI's chat completion API format.
std::vector<common_chat_msg> common_chat_msgs_parse_oaicompat(const nlohmann::ordered_json & messages);
std::vector<common_chat_msg> common_chat_msgs_parse_oaicompat(const common_json & messages);
std::vector<common_chat_tool> common_chat_tools_parse_oaicompat(const nlohmann::ordered_json & tools);
std::vector<common_chat_tool> common_chat_tools_parse_oaicompat(const common_json & tools);
common_chat_continuation common_chat_continuation_parse(const nlohmann::ordered_json & value);
common_chat_continuation common_chat_continuation_parse(const common_json & value);
// DEPRECATED: only used in tests
nlohmann::ordered_json common_chat_msgs_to_json_oaicompat(const std::vector<common_chat_msg> & msgs, bool concat_typed_text = false);
common_json common_chat_msgs_to_json_oaicompat(const std::vector<common_chat_msg> & msgs, bool concat_typed_text = false);
nlohmann::ordered_json common_chat_tools_to_json_oaicompat(const std::vector<common_chat_tool> & tools);
common_json common_chat_tools_to_json_oaicompat(const std::vector<common_chat_tool> & tools);
// get template caps, useful for reporting to server /props endpoint
std::map<std::string, bool> common_chat_templates_get_caps(const common_chat_templates * chat_templates);
@@ -386,4 +385,4 @@ struct common_chat_prompt_preset {
common_chat_prompt_preset common_chat_get_asr_prompt(const common_chat_templates * chat_templates);
common_chat_msg_delimiters common_chat_msg_delimiters_parse(const nlohmann::ordered_json & delimiters);
common_chat_msg_delimiters common_chat_msg_delimiters_parse(const common_json & delimiters);
common/common.cpp
+3 -2
View File
@@ -408,10 +408,11 @@ void common_params_print_info(const common_params & params, bool print_devices)
#endif
COM_TRC("%s: build %d (%s) with %s for %s%s\n", __func__, llama_build_number(), llama_commit(), llama_compiler(), llama_build_target(), build_type);
COM_INF("%s: verbosity = %d (adjust with the `-lv N` CLI arg)\n", __func__, common_log_get_verbosity_thold());
const int verbosity = common_log_get_verbosity_thold();
COM_INF("%s: verbosity = %d (adjust with the `-lv N` CLI arg)\n", __func__, verbosity);
// device enumeration creates a primary context on CUDA backends, skip it when the caller does not own any device
if (print_devices) {
if (print_devices && verbosity >= LOG_LEVEL_TRACE) {
COM_TRC("%s", "device_info:\n");
for (size_t i = 0; i < ggml_backend_dev_count(); ++i) {
auto * dev = ggml_backend_dev_get(i);
common/download.cpp
+6 -10
View File
@@ -5,9 +5,7 @@
#include "log.h"
#include "download.h"
#include "hf-cache.h"
#define JSON_ASSERT GGML_ASSERT
#include <nlohmann/json.hpp>
#include "json.h"
#include <algorithm>
#include <filesystem>
@@ -44,8 +42,6 @@
#include <unistd.h>
#endif
using json = nlohmann::ordered_json;
//
// downloader
//
@@ -856,8 +852,8 @@ static std::string common_docker_get_token(const std::string & repo) {
throw std::runtime_error("Failed to get Docker registry token, HTTP code: " + std::to_string(res.first));
}
std::string response_str(res.second.begin(), res.second.end());
nlohmann::ordered_json response = nlohmann::ordered_json::parse(response_str);
std::string response_str(res.second.begin(), res.second.end());
common_json response = common_json::parse(response_str);
if (!response.contains("token")) {
throw std::runtime_error("Docker registry token response missing 'token' field");
@@ -919,9 +915,9 @@ std::string common_docker_resolve_model(const std::string & docker) {
throw std::runtime_error("Failed to get Docker manifest, HTTP code: " + std::to_string(manifest_res.first));
}
std::string manifest_str(manifest_res.second.begin(), manifest_res.second.end());
nlohmann::ordered_json manifest = nlohmann::ordered_json::parse(manifest_str);
std::string gguf_digest; // Find the GGUF layer
std::string manifest_str(manifest_res.second.begin(), manifest_res.second.end());
common_json manifest = common_json::parse(manifest_str);
std::string gguf_digest; // Find the GGUF layer
if (manifest.contains("layers")) {
for (const auto & layer : manifest["layers"]) {
if (layer.contains("mediaType")) {
common/hf-cache.cpp
+7 -11
View File
@@ -4,9 +4,7 @@
#include "common.h"
#include "log.h"
#include "http.h"
#define JSON_ASSERT GGML_ASSERT
#include <nlohmann/json.hpp>
#include "json.h"
#include <filesystem>
#include <fstream>
@@ -15,8 +13,6 @@
#include <string_view>
#include <stdexcept>
namespace nl = nlohmann;
#if defined(_WIN32)
#define WIN32_LEAN_AND_MEAN
#ifndef NOMINMAX
@@ -195,8 +191,8 @@ static void safe_write_file(const fs::path & path, const std::string & data) {
}
}
static nl::json api_get(const std::string & url,
const std::string & token) {
static common_json api_get(const std::string & url,
const std::string & token) {
auto [cli, parts] = common_http_client(url);
httplib::Headers headers = {
@@ -214,10 +210,10 @@ static nl::json api_get(const std::string & url,
auto body = res->body;
if (res->status == 200) {
return nl::json::parse(res->body);
return common_json::parse(res->body);
}
try {
body = nl::json::parse(res->body)["error"].get<std::string>();
body = common_json::parse(res->body)["error"].get<std::string>();
} catch (...) { }
throw std::runtime_error("GET failed (" + std::to_string(res->status) + "): " + body);
@@ -280,7 +276,7 @@ static std::string get_repo_commit(const std::string & repo_id,
safe_write_file(refs_path / name, commit);
return commit;
} catch (const nl::json::exception & e) {
} catch (const common_json_error & e) {
LOG_ERR("%s: JSON error: %s\n", __func__, e.what());
} catch (const std::exception & e) {
LOG_ERR("%s: error: %s\n", __func__, e.what());
@@ -358,7 +354,7 @@ hf_files get_repo_files(const std::string & repo_id,
files.push_back(file);
}
} catch (const nl::json::exception & e) {
} catch (const common_json_error & e) {
LOG_ERR("%s: JSON error: %s\n", __func__, e.what());
} catch (const std::exception & e) {
LOG_ERR("%s: error: %s\n", __func__, e.what());
common/jinja/README.md
+1 -1
View File
@@ -7,7 +7,7 @@ The implementation can be found in the `common/jinja` directory.
## Key Features
- Input marking: security against special token injection
- Decoupled from `nlohmann::json`: this dependency is only used for JSON-to-internal type translation and is completely optional
- Decoupled from the JSON library: `common_json` is only used for JSON-to-internal type translation and is completely optional
- Minimal primitive types: int, float, bool, string, array, object, none, undefined
- Detailed logging: allow source tracing on error
- Clean architecture: workarounds are applied to input data before entering the runtime (see `common/chat.cpp`)
common/jinja/caps.cpp
+2 -2
View File
@@ -4,7 +4,7 @@
// note: the json dependency is only for defining input in a convenient way
// we can remove it in the future when we figure out a better way to define inputs using jinja::value
#include <nlohmann/json.hpp>
#include "json.h"
#include <functional>
#include <sstream>
@@ -14,7 +14,7 @@
#endif
#define FILENAME "jinja-caps"
using json = nlohmann::ordered_json;
using json = common_json;
namespace jinja {
common/jinja/value.cpp
+3 -3
View File
@@ -3,7 +3,7 @@
#include "value.h"
// for converting from JSON to jinja values
#include <nlohmann/json.hpp>
#include "json.h"
#include <sstream>
#include <string>
@@ -1358,7 +1358,7 @@ const func_builtins & value_undefined_t::get_builtins() const {
//////////////////////////////////
static value from_json(const nlohmann::ordered_json & j, bool mark_input) {
static value from_json(const common_json & j, bool mark_input) {
if (j.is_null()) {
return mk_val<value_none>();
} else if (j.is_boolean()) {
@@ -1455,7 +1455,7 @@ bool value_compare(const value & a, const value & b, value_compare_op op) {
}
template<>
void global_from_json(context & ctx, const nlohmann::ordered_json & json_obj, bool mark_input) {
void global_from_json(context & ctx, const common_json & json_obj, bool mark_input) {
// printf("global_from_json: %s\n" , json_obj.dump(2).c_str());
if (json_obj.is_null() || !json_obj.is_object()) {
throw std::runtime_error("global_from_json: input JSON value must be an object");
common/jinja/value.h
+1 -1
View File
@@ -86,7 +86,7 @@ struct context; // forward declaration
// marking input can be useful for tracking data provenance
// and preventing template injection attacks
//
// Note: T_JSON can be nlohmann::ordered_json
// Note: T_JSON can be common_json
template<typename T_JSON>
void global_from_json(context & ctx, const T_JSON & json_obj, bool mark_input);
common/json-schema-to-grammar.cpp
+12 -9
View File
@@ -1,9 +1,8 @@
#include "json-schema-to-grammar.h"
#include "common.h"
#include <nlohmann/json.hpp>
#include <algorithm>
#include <limits>
#include <map>
#include <regex>
#include <sstream>
@@ -12,7 +11,7 @@
#include <unordered_set>
#include <vector>
using json = nlohmann::ordered_json;
using json = common_json;
static std::string build_repetition(const std::string & item_rule, int min_items, int max_items, const std::string & separator_rule = "") {
auto has_max = max_items != std::numeric_limits<int>::max();
@@ -917,7 +916,11 @@ public:
return _add_rule(rule_name, _resolve_ref(schema["$ref"]));
}
if (schema.contains("oneOf") || schema.contains("anyOf")) {
std::vector<json> alt_schemas = schema.contains("oneOf") ? schema["oneOf"].get<std::vector<json>>() : schema["anyOf"].get<std::vector<json>>();
const json & alts = schema.contains("oneOf") ? schema.at("oneOf") : schema.at("anyOf");
std::vector<json> alt_schemas;
for (const auto & alt : alts) {
alt_schemas.push_back(alt);
}
return _add_rule(rule_name, _generate_union_rule(name, alt_schemas));
}
if (schema_type.is_array()) {
@@ -1111,7 +1114,7 @@ common_schema_info::~common_schema_info() = default;
common_schema_info::common_schema_info(common_schema_info &&) noexcept = default;
common_schema_info & common_schema_info::operator=(common_schema_info &&) noexcept = default;
void common_schema_info::resolve_refs(nlohmann::ordered_json & schema) {
void common_schema_info::resolve_refs(common_json & schema) {
impl_->resolve_refs(schema, "");
}
@@ -1119,7 +1122,7 @@ void common_schema_info::resolve_refs(nlohmann::ordered_json & schema) {
// Some models emit raw string values rather than JSON-encoded strings for string parameters.
// If any branch of the schema (via oneOf, anyOf, $ref, etc.) permits a string, this returns
// true, allowing callers to handle the value as a raw string for simplicity.
bool common_schema_info::resolves_to_string(const nlohmann::ordered_json & schema) {
bool common_schema_info::resolves_to_string(const common_json & schema) {
std::unordered_set<std::string> visited_refs;
std::function<bool(const json &)> check = [&](const json & s) -> bool {
@@ -1227,7 +1230,7 @@ bool common_schema_info::resolves_to_string(const nlohmann::ordered_json & schem
return check(schema);
}
std::string json_schema_to_grammar(const json & schema, bool force_gbnf) {
std::string json_schema_to_grammar(const common_json & schema, bool force_gbnf) {
#ifdef LLAMA_USE_LLGUIDANCE
if (!force_gbnf) {
return "%llguidance {}\nstart: %json " + schema.dump();
@@ -1248,10 +1251,10 @@ std::string build_grammar(const std::function<void(const common_grammar_builder
/* .add_rule = */ [&](const std::string & name, const std::string & rule) {
return converter._add_rule(name, rule);
},
/* .add_schema = */ [&](const std::string & name, const nlohmann::ordered_json & schema) {
/* .add_schema = */ [&](const std::string & name, const common_json & schema) {
return converter.visit(schema, name == "root" ? "" : name);
},
/* .resolve_refs = */ [&](nlohmann::ordered_json & schema) {
/* .resolve_refs = */ [&](common_json & schema) {
converter.resolve_refs(schema, "");
}
};
common/json-schema-to-grammar.h
+6 -6
View File
@@ -1,12 +1,12 @@
#pragma once
#include <nlohmann/json_fwd.hpp>
#include "json.h"
#include <functional>
#include <memory>
#include <string>
std::string json_schema_to_grammar(const nlohmann::ordered_json & schema,
std::string json_schema_to_grammar(const common_json & schema,
bool force_gbnf = false);
class common_schema_converter;
@@ -24,14 +24,14 @@ class common_schema_info {
common_schema_info(common_schema_info &&) noexcept;
common_schema_info & operator=(common_schema_info &&) noexcept;
void resolve_refs(nlohmann::ordered_json & schema);
bool resolves_to_string(const nlohmann::ordered_json & schema);
void resolve_refs(common_json & schema);
bool resolves_to_string(const common_json & schema);
};
struct common_grammar_builder {
std::function<std::string(const std::string &, const std::string &)> add_rule;
std::function<std::string(const std::string &, const nlohmann::ordered_json &)> add_schema;
std::function<void(nlohmann::ordered_json &)> resolve_refs;
std::function<std::string(const std::string &, const common_json &)> add_schema;
std::function<void(common_json &)> resolve_refs;
};
struct common_grammar_options {
common/json.cpp
+433
View File
@@ -0,0 +1,433 @@
#include "json.h"
#include "ggml.h"
#define JSON_ASSERT GGML_ASSERT
#include <nlohmann/json.hpp>
#include <iterator>
#include <new>
#include <set>
#include <unordered_map>
#include <vector>
using nlohmann::ordered_json;
// a common_json is the backing value, so any value of a tree can be used as a common_json
static_assert(sizeof(ordered_json) <= sizeof(common_json), "common_json storage is too small");
static_assert(alignof(ordered_json) <= alignof(common_json), "common_json alignment is too weak");
// runs fn and gives every error of the backing library as a common_json_error
template <typename F>
static decltype(auto) guard(F && fn) {
try {
return fn();
} catch (const ordered_json::exception & e) {
throw common_json_error(e.what());
}
}
static ordered_json & as_json(common_json * self) {
return *reinterpret_cast<ordered_json *>(self);
}
static const ordered_json & as_json(const common_json * self) {
return *reinterpret_cast<const ordered_json *>(self);
}
static common_json & as_common(ordered_json & json) {
return *reinterpret_cast<common_json *>(&json);
}
static const common_json & as_common(const ordered_json & json) {
return *reinterpret_cast<const common_json *>(&json);
}
static ordered_json to_json(const common_json_value & val) {
switch (val.type) {
case common_json_value::VAL_NULL: return nullptr;
case common_json_value::VAL_BOOL: return val.val_bool;
case common_json_value::VAL_INT: return val.val_int;
case common_json_value::VAL_UINT: return val.val_uint;
case common_json_value::VAL_DOUBLE: return val.val_double;
case common_json_value::VAL_STRING: return val.val_string;
case common_json_value::VAL_JSON:
// one owner means no one else can see this tree, so it is safe to move it out
// note: this makes a value single use, same as the json_ref of the backing library
if (val.val_json.use_count() == 1) {
return std::move(as_json(val.val_json.get()));
}
return as_json(val.val_json.get());
}
return nullptr;
}
common_json_value::common_json_value(const char * val) {
if (val) {
type = VAL_STRING;
val_string = val;
} else {
type = VAL_NULL;
}
}
common_json_value::common_json_value(const common_json & val) :
type(VAL_JSON), val_json(std::make_shared<common_json>(val)) {}
common_json_value::common_json_value(common_json && val) :
type(VAL_JSON), val_json(std::make_shared<common_json>(std::move(val))) {}
// the ctors and get<T>() below are explicit specializations, giving strong symbols
// an explicit instantiation is a weak symbol, dropped by some LTO builds (clang-cl)
template <typename T>
static std::shared_ptr<common_json> set_json(const std::set<T> & vals) {
common_json out = common_json::array();
for (const auto & val : vals) {
out.push_back(val);
}
return std::make_shared<common_json>(std::move(out));
}
// a set value is usable only for the types below
#define COMMON_JSON_SET(...) template <> common_json_value::common_json_value(const std::set<__VA_ARGS__> & vals) : type(VAL_JSON), val_json(set_json(vals)) {}
COMMON_JSON_SET(int)
COMMON_JSON_SET(std::string)
#undef COMMON_JSON_SET
template <typename T>
static std::shared_ptr<common_json> map_json(const T & vals) {
common_json out = common_json::object();
for (const auto & val : vals) {
out.set({ val.first, val.second });
}
return std::make_shared<common_json>(std::move(out));
}
// a map value is usable only for the types below
#define COMMON_JSON_MAP(...) template <> common_json_value::common_json_value(const std::map<std::string, __VA_ARGS__> & vals) : type(VAL_JSON), val_json(map_json(vals)) {}
COMMON_JSON_MAP(bool)
COMMON_JSON_MAP(std::string)
#undef COMMON_JSON_MAP
// an unordered map value is usable only for the types below
#define COMMON_JSON_UMAP(...) template <> common_json_value::common_json_value(const std::unordered_map<std::string, __VA_ARGS__> & vals) : type(VAL_JSON), val_json(map_json(vals)) {}
COMMON_JSON_UMAP(size_t)
#undef COMMON_JSON_UMAP
template <typename T>
static std::shared_ptr<common_json> vec_json(const std::vector<T> & vals) {
common_json out = common_json::array();
for (const auto & val : vals) {
out.push_back(val);
}
return std::make_shared<common_json>(std::move(out));
}
// a vector value is usable only for the types below
// note: std::vector<bool> is not here, its proxy reference does not convert
#define COMMON_JSON_VEC(...) template <> common_json_value::common_json_value(const std::vector<__VA_ARGS__> & vals) : type(VAL_JSON), val_json(vec_json(vals)) {}
COMMON_JSON_VEC(int)
COMMON_JSON_VEC(unsigned char)
COMMON_JSON_VEC(unsigned int)
COMMON_JSON_VEC(long)
COMMON_JSON_VEC(unsigned long)
COMMON_JSON_VEC(long long)
COMMON_JSON_VEC(unsigned long long)
COMMON_JSON_VEC(float)
COMMON_JSON_VEC(double)
COMMON_JSON_VEC(std::string)
COMMON_JSON_VEC(std::vector<float>)
COMMON_JSON_VEC(common_json)
#undef COMMON_JSON_VEC
common_json_value::common_json_value(std::initializer_list<common_json_item> items) :
type(VAL_JSON), val_json(std::make_shared<common_json>(items)) {}
// null, same as the backing library
// operator[] turns it into an object, push_back() into an array
common_json::common_json() {
new (storage) ordered_json();
}
common_json::common_json(const common_json & other) {
new (storage) ordered_json(as_json(&other));
}
common_json::common_json(common_json && other) noexcept {
new (storage) ordered_json(std::move(as_json(&other)));
}
common_json::common_json(std::initializer_list<common_json_item> items) {
new (storage) ordered_json(ordered_json::object());
for (const auto & item : items) {
set(item);
}
}
common_json::common_json(const common_json_value & val) {
new (storage) ordered_json(to_json(val));
}
common_json::common_json(std::nullptr_t) {
new (storage) ordered_json(nullptr);
}
common_json & common_json::operator=(common_json other) noexcept {
as_json(this).swap(as_json(&other));
return *this;
}
common_json::~common_json() {
as_json(this).~basic_json();
}
common_json common_json::parse(const std::string & text) {
try {
// the assignment moves the parsed tree in, it does not copy
common_json out;
as_json(&out) = ordered_json::parse(text);
return out;
} catch (const std::exception & e) {
throw common_json_error(e.what());
}
}
common_json common_json::parse_no_throw(const std::string & text) {
common_json out;
as_json(&out) = ordered_json::parse(text, nullptr, false);
return out;
}
bool common_json::is_discarded() const {
return as_json(this).is_discarded();
}
common_json common_json::array() {
common_json out;
as_json(&out) = ordered_json::array();
return out;
}
common_json common_json::array(std::initializer_list<common_json_value> vals) {
common_json out;
ordered_json & arr = as_json(&out);
arr = ordered_json::array();
for (const auto & val : vals) {
arr.push_back(to_json(val));
}
return out;
}
common_json common_json::object() {
common_json out;
as_json(&out) = ordered_json::object();
return out;
}
common_json common_json::object(std::initializer_list<common_json_item> items) {
return common_json(items);
}
common_json common_json::make(const common_json_value & val) {
return common_json(val);
}
bool common_json::is_null() const { return as_json(this).is_null(); }
bool common_json::is_object() const { return as_json(this).is_object(); }
bool common_json::is_array() const { return as_json(this).is_array(); }
bool common_json::is_string() const { return as_json(this).is_string(); }
bool common_json::is_boolean() const { return as_json(this).is_boolean(); }
bool common_json::is_number() const { return as_json(this).is_number(); }
bool common_json::is_number_integer() const { return as_json(this).is_number_integer(); }
bool common_json::is_number_float() const { return as_json(this).is_number_float(); }
bool common_json::empty() const { return as_json(this).empty(); }
size_t common_json::size() const { return as_json(this).size(); }
bool common_json::contains(const std::string & key) const {
return as_json(this).contains(key);
}
bool common_json::operator==(const common_json_value & val) const {
// compare a tree in place, to_json() would copy it
if (val.type == common_json_value::VAL_JSON) {
return as_json(this) == as_json(val.val_json.get());
}
return as_json(this) == to_json(val);
}
bool common_json::operator!=(const common_json_value & val) const {
return !(*this == val);
}
common_json & common_json::at(const std::string & key) { return guard([&]() -> common_json & { return as_common(as_json(this).at(key)); }); }
const common_json & common_json::at(const std::string & key) const { return guard([&]() -> const common_json & { return as_common(as_json(this).at(key)); }); }
common_json & common_json::at(size_t idx) { return guard([&]() -> common_json & { return as_common(as_json(this).at(idx)); }); }
const common_json & common_json::at(size_t idx) const { return guard([&]() -> const common_json & { return as_common(as_json(this).at(idx)); }); }
common_json & common_json::operator[](const std::string & key) { return guard([&]() -> common_json & { return as_common(as_json(this)[key]); }); }
const common_json & common_json::operator[](const std::string & key) const { return guard([&]() -> const common_json & { return as_common(as_json(this).at(key)); }); }
common_json & common_json::operator[](size_t idx) { return guard([&]() -> common_json & { return as_common(as_json(this)[idx]); }); }
const common_json & common_json::operator[](size_t idx) const { return guard([&]() -> const common_json & { return as_common(as_json(this).at(idx)); }); }
common_json & common_json::front() { return as_common(as_json(this).front()); }
const common_json & common_json::front() const { return as_common(as_json(this).front()); }
common_json & common_json::back() { return as_common(as_json(this).back()); }
const common_json & common_json::back() const { return as_common(as_json(this).back()); }
void common_json::clear() {
as_json(this).clear();
}
void common_json::erase(const std::string & key) {
guard([&] { as_json(this).erase(key); });
}
void common_json::erase(size_t idx) {
guard([&] { as_json(this).erase(idx); });
}
void common_json::assign(const common_json_value & val) {
as_json(this) = to_json(val);
}
void common_json::set(const common_json_item & item) {
guard([&] { as_json(this)[item.key] = to_json(item.val); });
}
void common_json::push_back(const common_json_value & val) {
guard([&] { as_json(this).push_back(to_json(val)); });
}
void common_json::push_back(std::initializer_list<common_json_item> items) {
common_json val(items);
guard([&] { as_json(this).push_back(std::move(as_json(&val))); });
}
size_t common_json::count(const std::string & key) const {
return as_json(this).count(key);
}
void common_json::insert(const common_json & vals) {
guard([&] {
ordered_json & self = as_json(this);
self.insert(self.end(), as_json(&vals).begin(), as_json(&vals).end());
});
}
std::string common_json::dump(int indent) const {
return guard([&] { return as_json(this).dump(indent); });
}
std::string common_json::dump_safe(int indent) const {
return as_json(this).dump(indent, ' ', false, ordered_json::error_handler_t::replace);
}
// an array is indexed directly, an object needs a walk from the start
common_json & common_json::iterator::operator*() const {
return guard([&]() -> common_json & {
ordered_json & j = as_json(node);
if (j.is_object()) {
return as_common(std::next(j.begin(), idx).value());
}
if (j.is_array()) {
return as_common(j[idx]);
}
// a plain value gives itself once, same as the backing library
return *node;
});
}
std::string common_json::iterator::key() const {
return guard([&] { return std::next(as_json(node).begin(), idx).key(); });
}
common_json::iterator common_json::begin() const {
return iterator(const_cast<common_json *>(this), 0);
}
common_json::iterator common_json::end() const {
return iterator(const_cast<common_json *>(this), size());
}
// the keys follow the backing library: the index for an array, "" for a plain value
common_json::items_view::entry common_json::items_view::iterator::operator*() const {
return guard([&]() -> entry {
ordered_json & j = as_json(node);
if (j.is_object()) {
auto it = std::next(j.begin(), idx);
return { it.key(), as_common(it.value()) };
}
if (j.is_array()) {
return { std::to_string(idx), as_common(j[idx]) };
}
return { std::string(), *node };
});
}
common_json::items_view common_json::items() const {
return items_view(const_cast<common_json *>(this), size());
}
// the backing library cannot build a common_json, so this one is just a copy
template <> common_json common_json::get<common_json>() const {
return *this;
}
// get<T>() is usable only for the types below
#define COMMON_JSON_GET(...) template <> __VA_ARGS__ common_json::get<__VA_ARGS__>() const { return guard([&] { return as_json(this).get<__VA_ARGS__>(); }); }
COMMON_JSON_GET(bool)
COMMON_JSON_GET(int)
COMMON_JSON_GET(unsigned int)
COMMON_JSON_GET(long)
COMMON_JSON_GET(unsigned long)
COMMON_JSON_GET(long long)
COMMON_JSON_GET(unsigned long long)
COMMON_JSON_GET(float)
COMMON_JSON_GET(double)
COMMON_JSON_GET(std::string)
COMMON_JSON_GET(std::vector<float>)
COMMON_JSON_GET(std::vector<std::string>)
COMMON_JSON_GET(std::set<std::string>)
COMMON_JSON_GET(std::vector<int>)
COMMON_JSON_GET(std::vector<size_t>)
COMMON_JSON_GET(std::unordered_map<std::string, size_t>)
#undef COMMON_JSON_GET
// must stay below the get<std::string> specialization
common_json::operator std::string() const {
return get<std::string>();
}
std::string common_json::value(const std::string & key, const char * def) const {
return contains(key) ? at(key).get<std::string>() : std::string(def);
}
common/json.h
+352
View File
@@ -0,0 +1,352 @@
#pragma once
#include <cstddef>
#include <cstdint>
#include <initializer_list>
#include <iterator>
#include <map>
#include <memory>
#include <set>
#include <stdexcept>
#include <string>
#include <string_view>
#include <type_traits>
#include <unordered_map>
#include <utility>
#include <vector>
// common_json, a thin wrapper around vendor json library
// the underlay library is pimpl, we are using nlohmann::json for now
//
// many features of the library are deliberately left out, to keep this interface small and generic and to keep compile time down
//
// some main differences compared to nlohmann::json :
// - object keys keep the order in which they are added
// - errors are always throw as common_json_error
// - obj.push_back({key, val}) is intentionally unsupported to avoid confusion with push_back on a vector; write it as obj[key] = val for clarity
// - a braced pair in value position does not build, e.g. {"key", {"a", "b"}}; write array({"a", "b"}) where nlohmann made an array
//
// in doubt, search the code base for an existing usage example; do not add anything to this header unless absolutely necessary
class common_json;
// common_json_value holds a list of these, and each of them holds a value, so one must come first
struct common_json_item;
struct common_json_error : std::runtime_error {
using std::runtime_error::runtime_error;
};
// one value, tagged so that this header stays free of the backing library
// note: a value that holds a tree is single use, the second use gives null
struct common_json_value {
enum value_type {
VAL_NULL,
VAL_BOOL,
VAL_INT,
VAL_UINT,
VAL_DOUBLE,
VAL_STRING,
VAL_JSON,
};
value_type type = VAL_NULL;
union {
bool val_bool;
int64_t val_int;
uint64_t val_uint = 0;
double val_double;
};
std::string val_string;
std::shared_ptr<common_json> val_json;
common_json_value(std::nullptr_t = nullptr) : type(VAL_NULL) {}
common_json_value(bool val) : type(VAL_BOOL), val_bool(val) {}
common_json_value(std::string val) : type(VAL_STRING), val_string(std::move(val)) {}
// without this a string_view lands on the common_json ctor below and recurses
common_json_value(std::string_view val) : type(VAL_STRING), val_string(val) {}
common_json_value(const char * val);
common_json_value(const common_json & val);
common_json_value(common_json && val);
// only for the types instantiated in json.cpp, the rest fails at link time
template <typename T> common_json_value(const std::vector<T> & vals);
// a set becomes an array, in the set's own order
template <typename T> common_json_value(const std::set<T> & vals);
// a map becomes an object, keyed in the map's own order
template <typename T> common_json_value(const std::map<std::string, T> & vals);
template <typename T> common_json_value(const std::unordered_map<std::string, T> & vals);
// nested object, e.g. {"fn", {{"name", "x"}}}
// note: a nested pair {"a", "b"} does not build, use common_json::array({"a", "b"}) for an array
common_json_value(std::initializer_list<common_json_item> items);
template <typename T, typename std::enable_if<std::is_integral<T>::value && !std::is_same<T, bool>::value, int>::type = 0>
common_json_value(T val) : type(std::is_signed<T>::value ? VAL_INT : VAL_UINT) {
if (std::is_signed<T>::value) {
val_int = (int64_t) val;
} else {
val_uint = (uint64_t) val;
}
}
template <typename T, typename std::enable_if<std::is_floating_point<T>::value, int>::type = 0>
common_json_value(T val) : type(VAL_DOUBLE), val_double((double) val) {}
};
struct common_json_item {
std::string key;
common_json_value val;
template <typename T>
common_json_item(std::string key, T && val) :
key(std::move(key)), val(std::forward<T>(val)) {}
// a braced list cannot deduce T, so it needs its own overload
common_json_item(std::string key, std::initializer_list<common_json_item> items) :
key(std::move(key)), val(items) {}
};
// the types common_json_value holds on its own
// anything else reaches its common_json ctor and recurses forever
template <typename T> struct common_json_is_value : std::integral_constant<bool,
std::is_arithmetic<T>::value ||
std::is_same<T, std::nullptr_t>::value ||
std::is_same<T, std::string>::value ||
std::is_same<T, std::string_view>::value ||
std::is_same<T, char *>::value ||
std::is_same<T, const char *>::value ||
std::is_same<T, common_json>::value> {};
template <typename T, typename A>
struct common_json_is_value<std::vector<T, A>> : std::true_type {};
template <typename T, typename C, typename A>
struct common_json_is_value<std::set<T, C, A>> : std::true_type {};
template <typename V, typename C, typename A>
struct common_json_is_value<std::map<std::string, V, C, A>> : std::true_type {};
template <typename V, typename H, typename E, typename A>
struct common_json_is_value<std::unordered_map<std::string, V, H, E, A>> : std::true_type {};
class common_json {
public:
common_json();
common_json(const common_json & other);
common_json(common_json && other) noexcept;
common_json(std::initializer_list<common_json_item> items);
common_json(const common_json_value & val);
// direct, a value would need two conversions in a row
common_json(std::nullptr_t);
// one step, so that "abc" or a vector can go straight into a common_json
template <typename T, typename std::enable_if<!std::is_same<typename std::decay<T>::type, common_json>::value &&
!std::is_same<typename std::decay<T>::type, common_json_value>::value, int>::type = 0>
common_json(T && val) : common_json(common_json_value(std::forward<T>(val))) {
static_assert(common_json_is_value<typename std::decay<T>::type>::value,
"no common_json_value ctor holds this type, add one instead of letting it recurse");
}
// by value, same as the backing library
// the right side is copied before the left side can invalidate it, e.g. msg["a"] = msg.at("b")
common_json & operator=(common_json other) noexcept;
~common_json();
// throws common_json_error if the text is not valid JSON
static common_json parse(const std::string & text);
// gives a discarded value instead of throwing, check it with is_discarded()
static common_json parse_no_throw(const std::string & text);
bool is_discarded() const;
static common_json array();
static common_json array(std::initializer_list<common_json_value> vals);
static common_json object();
static common_json object(std::initializer_list<common_json_item> items);
// holds a single value, e.g. make("abc").dump() gives "\"abc\""
static common_json make(const common_json_value & val);
bool is_null() const;
bool is_object() const;
bool is_array() const;
bool is_string() const;
bool is_boolean() const;
bool is_number() const;
bool is_number_integer() const;
bool is_number_float() const;
bool empty() const;
size_t size() const;
bool contains(const std::string & key) const;
bool operator==(const common_json_value & val) const;
bool operator!=(const common_json_value & val) const;
// at() throws common_json_error if the key is missing, operator[] adds a null value instead
// note: a const operator[] cannot add, it throws like at()
common_json & at(const std::string & key);
const common_json & at(const std::string & key) const;
common_json & at(size_t idx);
const common_json & at(size_t idx) const;
common_json & operator[](const std::string & key);
const common_json & operator[](const std::string & key) const;
common_json & operator[](const char * key) { return (*this)[std::string(key)]; }
const common_json & operator[](const char * key) const { return (*this)[std::string(key)]; }
common_json & operator[](int idx) { return (*this)[to_idx(idx)]; }
const common_json & operator[](int idx) const { return (*this)[to_idx(idx)]; }
common_json & operator[](size_t idx);
const common_json & operator[](size_t idx) const;
common_json & front();
const common_json & front() const;
common_json & back();
const common_json & back() const;
void clear();
void erase(const std::string & key);
void erase(size_t idx);
// only for the types instantiated in json.cpp, the rest fails at link time
template <typename T> T get() const;
// implicit get<T>() for plain values, so they can be assigned to their C++ type directly
// note: kept to this short list on purpose, a wider one makes j["key"] ambiguous
// note: a numeric one would make "str = json;" ambiguous, a number converts to char too
operator std::string() const;
template <typename T>
T value(const std::string & key, T def) const {
return contains(key) ? at(key).get<T>() : def;
}
std::string value(const std::string & key, const char * def) const;
// a JSON default needs no get<T>(), it is already the right type
common_json value(const std::string & key, const common_json & def) const {
return contains(key) ? at(key) : def;
}
void assign(const common_json_value & val);
void set(const common_json_item & item);
void push_back(const common_json_value & val);
// appends one object, e.g. push_back({{"a", 1}})
void push_back(std::initializer_list<common_json_item> items);
// 1 if the key is there, 0 if not
size_t count(const std::string & key) const;
// appends every value of another array; inserting an array into itself throws
void insert(const common_json & vals);
// a common_json goes through the copy assignment above, everything else becomes a value
template <typename T, typename std::enable_if<!std::is_same<typename std::decay<T>::type, common_json>::value, int>::type = 0>
common_json & operator=(T && val) {
assign(common_json_value(std::forward<T>(val)));
return *this;
}
std::string dump(int indent = -1) const;
// same as dump(), but bad UTF-8 gets replaced instead of throwing
std::string dump_safe(int indent = -1) const;
// walks an array by index, or an object in insertion order
// a plain value gives itself once, same as the backing library
class iterator {
public:
using iterator_category = std::forward_iterator_tag;
using value_type = common_json;
using difference_type = std::ptrdiff_t;
using pointer = common_json *;
using reference = common_json &;
iterator(common_json * node, size_t idx) : node(node), idx(idx) {}
common_json & operator*() const;
common_json & value() const { return **this; }
std::string key() const;
iterator & operator++() {
idx++;
return *this;
}
bool operator!=(const iterator & other) const { return idx != other.idx; }
bool operator==(const iterator & other) const { return idx == other.idx; }
private:
common_json * node;
size_t idx;
};
iterator begin() const;
iterator end() const;
// allows: for (const auto & [key, val] : obj.items())
class items_view {
public:
// the members are public, so an entry also works with structured bindings
struct entry {
std::string k;
common_json & v;
const std::string & key() const { return k; }
common_json & value() const { return v; }
};
items_view(common_json * node, size_t n) : node(node), n(n) {}
class iterator {
public:
iterator(common_json * node, size_t idx) : node(node), idx(idx) {}
entry operator*() const;
iterator & operator++() {
idx++;
return *this;
}
bool operator!=(const iterator & other) const { return idx != other.idx; }
private:
common_json * node;
size_t idx;
};
iterator begin() const { return iterator(node, 0); }
iterator end() const { return iterator(node, n); }
private:
common_json * node;
size_t n;
};
items_view items() const;
private:
// a negative index must not turn into a huge size_t
static size_t to_idx(int idx) {
if (idx < 0) {
throw common_json_error("negative array index");
}
return (size_t) idx;
}
// the backing value is built here, json.cpp checks that it fits
// it cannot be a pointer: a value inside a tree would then not be a common_json
// at() could then only give back a copy instead of a real reference
alignas(8) unsigned char storage[32];
};
using common_json_entry = common_json::items_view::entry;
common/peg-parser.cpp
+15 -16
View File
@@ -10,7 +10,6 @@
#include <initializer_list>
#include <map>
#include <memory>
#include <nlohmann/json.hpp>
#include <regex>
#include <set>
#include <stdexcept>
@@ -1120,8 +1119,8 @@ common_peg_parser common_peg_parser_builder::chars(const std::string & classes,
return wrap(arena_.add_parser(common_peg_chars_parser{classes, ranges, negated, min, max}));
}
common_peg_parser common_peg_parser_builder::schema(const common_peg_parser & p, const std::string & name, const nlohmann::ordered_json & schema, bool raw) {
return wrap(arena_.add_parser(common_peg_schema_parser{p.id(), name, std::make_shared<nlohmann::ordered_json>(schema), raw}));
common_peg_parser common_peg_parser_builder::schema(const common_peg_parser & p, const std::string & name, const common_json & schema, bool raw) {
return wrap(arena_.add_parser(common_peg_schema_parser{p.id(), name, std::make_shared<common_json>(schema), raw}));
}
common_peg_parser common_peg_parser_builder::rule(const std::string & name, const common_peg_parser & p, bool trigger) {
@@ -1805,8 +1804,8 @@ void common_peg_arena::build_grammar(const common_grammar_builder & builder, boo
}
}
static nlohmann::json serialize_parser_variant(const common_peg_parser_variant & variant) {
using json = nlohmann::json;
static common_json serialize_parser_variant(const common_peg_parser_variant & variant) {
using json = common_json;
return std::visit([](const auto & p) -> json {
using T = std::decay_t<decltype(p)>;
@@ -1860,7 +1859,7 @@ static nlohmann::json serialize_parser_variant(const common_peg_parser_variant &
{"type", "schema"},
{"child", p.child},
{"name", p.name},
{"schema", p.schema ? *p.schema : nullptr},
{"schema", p.schema ? *p.schema : json(nullptr)},
{"raw", p.raw}
};
} else if constexpr (std::is_same_v<T, common_peg_rule_parser>) {
@@ -1888,19 +1887,19 @@ static nlohmann::json serialize_parser_variant(const common_peg_parser_variant &
}, variant);
}
nlohmann::json common_peg_arena::to_json() const {
auto parsers = nlohmann::json::array();
common_json common_peg_arena::to_json() const {
auto parsers = common_json::array();
for (const auto & parser : parsers_) {
parsers.push_back(serialize_parser_variant(parser));
}
return nlohmann::json{
return common_json{
{"parsers", parsers},
{"rules", rules_},
{"root", root_}
};
}
static common_peg_parser_variant deserialize_parser_variant(const nlohmann::json & j) {
static common_peg_parser_variant deserialize_parser_variant(const common_json & j) {
if (!j.contains("type") || !j["type"].is_string()) {
throw std::runtime_error("Parser variant JSON missing or invalid 'type' field");
}
@@ -1969,9 +1968,9 @@ static common_peg_parser_variant deserialize_parser_variant(const nlohmann::json
}
common_peg_chars_parser parser;
parser.pattern = j["pattern"];
parser.negated = j["negated"];
parser.min_count = j["min_count"];
parser.max_count = j["max_count"];
parser.negated = j["negated"].get<bool>();
parser.min_count = j["min_count"].get<int>();
parser.max_count = j["max_count"].get<int>();
for (const auto & range_json : j["ranges"]) {
if (!range_json.contains("start") || !range_json.contains("end")) {
throw std::runtime_error("char_range missing 'start' or 'end' field");
@@ -2007,7 +2006,7 @@ static common_peg_parser_variant deserialize_parser_variant(const nlohmann::json
parser.child = j["child"].get<common_peg_parser_id>();
parser.name = j["name"];
if (!j["schema"].is_null()) {
parser.schema = std::make_shared<nlohmann::ordered_json>(j["schema"]);
parser.schema = std::make_shared<common_json>(j["schema"]);
}
parser.raw = j["raw"].get<bool>();
return parser;
@@ -2069,7 +2068,7 @@ static common_peg_parser_variant deserialize_parser_variant(const nlohmann::json
throw std::runtime_error("Unknown parser type: " + type);
}
common_peg_arena common_peg_arena::from_json(const nlohmann::json & j) {
common_peg_arena common_peg_arena::from_json(const common_json & j) {
if (!j.contains("parsers") || !j["parsers"].is_array()) {
throw std::runtime_error("JSON missing or invalid 'parsers' array");
}
@@ -2109,7 +2108,7 @@ std::string common_peg_arena::save() const {
}
void common_peg_arena::load(const std::string & data) {
*this = from_json(nlohmann::json::parse(data));
*this = from_json(common_json::parse(data));
}
common_peg_arena build_peg_parser(const std::function<common_peg_parser(common_peg_parser_builder & builder)> & fn) {
common/peg-parser.h
+5 -5
View File
@@ -1,6 +1,6 @@
#pragma once
#include <nlohmann/json_fwd.hpp>
#include "json.h"
#include <memory>
#include <set>
@@ -245,7 +245,7 @@ struct common_peg_until_parser {
struct common_peg_schema_parser {
common_peg_parser_id child;
std::string name;
std::shared_ptr<nlohmann::ordered_json> schema;
std::shared_ptr<common_json> schema;
// Indicates if the GBNF should accept a raw string that matches the schema.
bool raw;
@@ -332,8 +332,8 @@ class common_peg_arena {
std::string dump(common_peg_parser_id id) const;
nlohmann::json to_json() const;
static common_peg_arena from_json(const nlohmann::json & j);
common_json to_json() const;
static common_peg_arena from_json(const common_json & j);
std::string save() const;
void load(const std::string & data);
@@ -490,7 +490,7 @@ class common_peg_parser_builder {
// Wraps a parser with JSON schema metadata for grammar generation.
// Used internally to convert JSON schemas to GBNF grammar rules.
common_peg_parser schema(const common_peg_parser & p, const std::string & name, const nlohmann::ordered_json & schema, bool raw = false);
common_peg_parser schema(const common_peg_parser & p, const std::string & name, const common_json & schema, bool raw = false);
// Creates a named rule, stores it in the grammar, and returns a ref.
// If trigger=true, marks this rule as an entry point for lazy grammar generation.
conversion/glm.py
+44 -5
View File
@@ -112,12 +112,38 @@ class GlmOCRModel(Glm4Model):
@ModelBase.example("zai-org/GLM-4.5-Air")
class Glm4MoeModel(TextModel):
model_arch = gguf.MODEL_ARCH.GLM4_MOE
supports_mtp_export = True
_n_main_layers: int | None = None
def __init__(self, *args, **kwargs):
super().__init__(*args, **kwargs)
# GLM4_MOE has num_hidden_layers + 1 actual layers (including NextN layer)
self.block_count = self.hparams["num_hidden_layers"] + self.hparams.get("num_nextn_predict_layers", 0)
self.tensor_map = gguf.get_tensor_name_map(self.model_arch, self.block_count)
if not self.no_mtp:
self.block_count += self.hparams.get("num_nextn_predict_layers", 0)
self.tensor_map = gguf.get_tensor_name_map(self.model_arch, self.block_count)
def index_tensors(self, remote_hf_model_id: str | None = None):
hparams = {**self.hparams, **self.hparams.get("text_config", {})}
key = next((k for k in ["n_layers", "num_hidden_layers", "n_layer", "num_layers"] if k in hparams), None)
type(self)._n_main_layers = hparams.get(key)
return super().index_tensors(remote_hf_model_id=remote_hf_model_id)
@classmethod
def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
if (titem := super().filter_tensors(item)) is None:
return None
name, gen = titem
assert cls._n_main_layers is not None
is_mtp = (m := re.match(r"model\.layers\.(\d+)\.", name)) is not None and int(m.group(1)) >= cls._n_main_layers
if is_mtp and cls.no_mtp:
return None
if cls.mtp_only and not is_mtp and name not in (
"model.embed_tokens.weight", "model.norm.weight", "lm_head.weight",
):
return None
return name, gen
def set_vocab(self):
return self._set_vocab_glm()
@@ -153,10 +179,22 @@ class Glm4MoeModel(TextModel):
if (norm_topk_prob := self.hparams.get("norm_topk_prob")) is not None:
self.gguf_writer.add_expert_weights_norm(norm_topk_prob)
# NextN/MTP prediction layers
if (num_nextn_predict_layers := self.hparams.get("num_nextn_predict_layers")) is not None:
if not self.no_mtp and (num_nextn_predict_layers := self.hparams.get("num_nextn_predict_layers")) is not None:
self.gguf_writer.add_nextn_predict_layers(num_nextn_predict_layers)
def prepare_metadata(self, vocab_only: bool):
from_dir = self.fname_out.is_dir()
super().prepare_metadata(vocab_only=vocab_only)
if not self.mtp_only or not from_dir:
return
output_type: str = self.ftype.name.partition("_")[2]
fname_default: str = gguf.naming_convention(
self.metadata.name, self.metadata.basename, self.metadata.finetune,
self.metadata.version, size_label=None, output_type=output_type, model_type=None)
self.fname_out = self.fname_out.parent / f"mtp-{fname_default}.gguf"
_experts: list[dict[str, Tensor]] | None = None
# note: unlike GLM4V non-MoE, we don't need to permute Q/K here since GLM4V_MOE uses Neox ordering already
@@ -348,6 +386,7 @@ class GlmMoeDsaModel(DeepseekV2Model):
@ModelBase.example("upstage/Solar-Open-100B")
class SolarOpenModel(Glm4MoeModel):
model_arch = gguf.MODEL_ARCH.GLM4_MOE
supports_mtp_export = False
def set_vocab(self):
from transformers import AutoTokenizer
ggml/include/ggml.h
+13 -8
View File
@@ -1742,6 +1742,19 @@ extern "C" {
struct ggml_tensor * a,
int n_past);
GGML_API struct ggml_tensor * ggml_clamp(
struct ggml_context * ctx,
struct ggml_tensor * a,
float min,
float max);
// in-place, returns view(a)
GGML_API struct ggml_tensor * ggml_clamp_inplace(
struct ggml_context * ctx,
struct ggml_tensor * a,
float min,
float max);
GGML_API struct ggml_tensor * ggml_soft_max(
struct ggml_context * ctx,
struct ggml_tensor * a);
@@ -2008,14 +2021,6 @@ extern "C" {
struct ggml_tensor * a,
int n_offs);
// clamp
// in-place, returns view(a)
GGML_API struct ggml_tensor * ggml_clamp(
struct ggml_context * ctx,
struct ggml_tensor * a,
float min,
float max);
// im2col
// converts data into a format that effectively results in a convolution when combined with matrix multiplication
GGML_API struct ggml_tensor * ggml_im2col(
ggml/src/ggml-alloc.c
+1
View File
@@ -40,6 +40,7 @@ bool ggml_op_can_inplace(enum ggml_op op) {
case GGML_OP_SILU_BACK:
case GGML_OP_RMS_NORM:
case GGML_OP_RMS_NORM_BACK:
case GGML_OP_CLAMP:
case GGML_OP_SOFT_MAX:
case GGML_OP_SOFT_MAX_BACK:
return true;
ggml/src/ggml-backend-meta.cpp
+251 -29
View File
@@ -592,7 +592,18 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
GGML_ASSERT(split_states_equal(src_ss[0], src_ss[1]));
return {assume_sync ? GGML_BACKEND_SPLIT_AXIS_MIRRORED : GGML_BACKEND_SPLIT_AXIS_PARTIAL, {0}, {1}, 1};
}
GGML_ABORT("fatal error");
if (src_ss[0].axis == src_ss[1].axis && src_ss[0].axis >= GGML_BACKEND_SPLIT_AXIS_2 &&
src_ss[0].axis < GGML_MAX_DIMS) {
GGML_ASSERT(split_states_equal(src_ss[0], src_ss[1]));
return src_ss[0];
}
// batched matmul with the batches split across devices and a replicated activation
if (src_ss[0].axis >= GGML_BACKEND_SPLIT_AXIS_2 && src_ss[0].axis < GGML_MAX_DIMS &&
src_ss[1].axis == GGML_BACKEND_SPLIT_AXIS_MIRRORED) {
return src_ss[0];
}
GGML_ABORT("unsupported mul_mat split states: node=%s src0=%s axis=%d src1=%s axis=%d",
tensor->name, tensor->src[0]->name, (int) src_ss[0].axis, tensor->src[1]->name, (int) src_ss[1].axis);
//return {GGML_BACKEND_SPLIT_AXIS_UNKNOWN, {0}, {1}, 1};
};
@@ -602,27 +613,40 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
case GGML_BACKEND_SPLIT_AXIS_1:
case GGML_BACKEND_SPLIT_AXIS_2:
case GGML_BACKEND_SPLIT_AXIS_3: {
GGML_ASSERT(src_ss[0].n_segments == 1);
if (src_ss[0].axis == ggml_n_dims(tensor->src[0]) - 1 && src_ss[0].nr[0] == 1) {
return {ggml_backend_meta_split_axis(ggml_n_dims(tensor) - 1), {0}, {1}, 1};
}
int64_t base_ne_in = tensor->src[0]->ne[0];
for (int dim = 1; dim <= src_ss[0].axis; dim++) {
int64_t base_ne_in = 1;
for (int dim = 0; dim <= src_ss[0].axis; dim++) {
base_ne_in *= tensor->src[0]->ne[dim];
}
base_ne_in /= src_ss[0].nr[0];
if (src_ss[0].n_segments == 1) {
base_ne_in /= src_ss[0].nr[0];
if (src_ss[0].axis == ggml_n_dims(tensor->src[0]) - 1 && src_ss[0].nr[0] == 1) {
return {ggml_backend_meta_split_axis(ggml_n_dims(tensor) - 1), {0}, {1}, 1};
}
if (src_ss[0].axis == GGML_BACKEND_SPLIT_AXIS_0 && tensor->ne[0] == tensor->src[0]->ne[0] &&
tensor->ne[1] == 1 && src_ss[0].nr[0] == 1) {
bool complete_rows = true;
for (size_t j = 0; j < n_bufs; j++) {
const int64_t ne = src_ss[0].ne[j];
complete_rows = complete_rows && (ne == 0 || ne == tensor->src[0]->ne[0]);
}
if (complete_rows) {
// Move a complete dim-0 split to the following singleton dimension.
return {GGML_BACKEND_SPLIT_AXIS_1, {0}, {1}, 1};
}
}
}
// Reshape outputs use one segment; split-state propagation merges source segments.
int64_t base_ne_out = 1;
for (int dim = 0; dim < GGML_MAX_DIMS; dim++) {
const int64_t base_ne_out_next = base_ne_out *= tensor->ne[dim];
if (base_ne_out_next % base_ne_in == 0) {
return {ggml_backend_meta_split_axis(dim), {0}, {uint32_t(base_ne_out_next/base_ne_in)}, 1};
base_ne_out *= tensor->ne[dim];
if (base_ne_out % base_ne_in == 0) {
return {ggml_backend_meta_split_axis(dim), {0}, {uint32_t(base_ne_out/base_ne_in)}, 1};
}
if (base_ne_out_next > base_ne_in) {
if (base_ne_out > base_ne_in) {
GGML_ASSERT(src_ss[0].n_segments == 1);
GGML_ASSERT(src_ss[0].nr[0] == 1);
return {ggml_backend_meta_split_axis(dim), {0}, {1}, 1};
}
base_ne_out = base_ne_out_next;
}
GGML_ABORT("shape mismatch for %s", ggml_op_name(tensor->op));
}
@@ -747,14 +771,33 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
};
auto handle_flash_attn_ext = [&](const std::vector<ggml_backend_meta_split_state> & src_ss) -> ggml_backend_meta_split_state {
GGML_ASSERT( src_ss[0].axis == GGML_BACKEND_SPLIT_AXIS_2);
GGML_ASSERT( src_ss[1].axis == GGML_BACKEND_SPLIT_AXIS_2);
GGML_ASSERT( src_ss[2].axis == GGML_BACKEND_SPLIT_AXIS_2);
GGML_ASSERT(tensor->src[4] == nullptr || src_ss[3].axis == GGML_BACKEND_SPLIT_AXIS_MIRRORED);
GGML_ASSERT(tensor->src[3] == nullptr || src_ss[3].axis == GGML_BACKEND_SPLIT_AXIS_MIRRORED);
if (src_ss[0].axis == GGML_BACKEND_SPLIT_AXIS_MIRRORED) {
GGML_ASSERT(src_ss[1].axis == GGML_BACKEND_SPLIT_AXIS_MIRRORED);
GGML_ASSERT(src_ss[2].axis == GGML_BACKEND_SPLIT_AXIS_MIRRORED);
GGML_ASSERT(tensor->src[4] == nullptr || src_ss[4].axis == GGML_BACKEND_SPLIT_AXIS_MIRRORED);
return {GGML_BACKEND_SPLIT_AXIS_MIRRORED, {0}, {1}, 1};
}
GGML_ASSERT(src_ss[0].axis == GGML_BACKEND_SPLIT_AXIS_2);
const bool kv_split = src_ss[1].axis == GGML_BACKEND_SPLIT_AXIS_2 &&
src_ss[2].axis == GGML_BACKEND_SPLIT_AXIS_2;
const bool kv_mirrored = src_ss[1].axis == GGML_BACKEND_SPLIT_AXIS_MIRRORED &&
src_ss[2].axis == GGML_BACKEND_SPLIT_AXIS_MIRRORED;
GGML_ASSERT(kv_split || kv_mirrored);
GGML_ASSERT(tensor->src[4] == nullptr || src_ss[4].axis == GGML_BACKEND_SPLIT_AXIS_0);
return {GGML_BACKEND_SPLIT_AXIS_1, {0}, {1}, 1};
};
auto handle_lightning_indexer = [&](
const std::vector<ggml_backend_meta_split_state> & src_ss) -> ggml_backend_meta_split_state {
for (size_t i = 0; i < 4; i++) {
GGML_ASSERT(src_ss[i].axis == GGML_BACKEND_SPLIT_AXIS_MIRRORED);
}
return {GGML_BACKEND_SPLIT_AXIS_MIRRORED, {0}, {1}, 1};
};
auto handle_ssm_conv = [&](const std::vector<ggml_backend_meta_split_state> & src_ss) -> ggml_backend_meta_split_state {
if (src_ss[0].axis == src_ss[1].axis) {
if (src_ss[0].axis == GGML_BACKEND_SPLIT_AXIS_0) {
@@ -792,7 +835,7 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
ggml_backend_dev_t dev = ggml_backend_buft_get_device(ggml_backend_buffer_get_type(tensor->buffer));
const ggml_backend_meta_device_context * dev_ctx = (const ggml_backend_meta_device_context *) dev->context;
ggml_backend_meta_split_state ret = dev_ctx->get_split_state(tensor, dev_ctx->get_split_state_ud);
if (ret.axis >= 0 && ret.axis <= GGML_MAX_DIMS) {
if (ret.axis >= 0 && ret.axis < GGML_MAX_DIMS) {
const int64_t granularity = ret.axis == GGML_BACKEND_SPLIT_AXIS_0 ? ggml_blck_size(tensor->type) : 1;
int64_t ne_sum = 0;
for (size_t s = 0; s < ret.n_segments; s++) {
@@ -802,6 +845,9 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
}
}
GGML_ASSERT(ne_sum == tensor->ne[ret.axis]);
} else if (ret.axis == GGML_BACKEND_SPLIT_AXIS_PARTIAL) {
GGML_ASSERT(ret.n_segments == 1);
GGML_ASSERT(ret.nr[0] == 1);
}
return ret;
}
@@ -922,7 +968,7 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
split_state = handle_rope(src_ss);
} break;
case GGML_OP_ROPE_BACK: {
split_state = handle_generic(src_ss, /*scalar_only =*/ true);
split_state = handle_rope(src_ss);
} break;
case GGML_OP_CLAMP: {
split_state = handle_generic(src_ss, /*scalar_only =*/ false);
@@ -986,6 +1032,9 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
case GGML_OP_GATED_DELTA_NET: {
split_state = handle_gated_delta_net(src_ss);
} break;
case GGML_OP_LIGHTNING_INDEXER: {
split_state = handle_lightning_indexer(src_ss);
} break;
case GGML_OP_DSV4_HC_COMB:
case GGML_OP_DSV4_HC_PRE:
case GGML_OP_DSV4_HC_POST: {
@@ -1070,13 +1119,14 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
if (buf_ctx->debug > 0) {
std::string srcs_info;
for (size_t i = 0; i < GGML_MAX_SRC; i++) {
if (tensor->src[i] == nullptr) {
if (tensor->src[i] == nullptr || tensor->src[i] == tensor) {
continue;
}
if (!srcs_info.empty()) {
srcs_info += ", ";
}
const ggml_backend_meta_split_state split_state = ggml_backend_meta_get_split_state(tensor->src[0], true);
const ggml_backend_meta_split_state split_state =
ggml_backend_meta_get_split_state(tensor->src[i], true);
GGML_ASSERT(split_state.n_segments == 1);
const char * axis_name = ggml_backend_meta_split_axis_name(split_state.axis);
std::string ne_info;
@@ -1255,6 +1305,108 @@ static enum ggml_status ggml_backend_meta_buffer_init_tensor(ggml_backend_buffer
return ggml_backend_meta_buffer_init_tensor_impl(buf_ctx->get_simple_tensor_container(tensor), tensor);
}
static void ggml_backend_meta_buffer_memset_tensor(
ggml_backend_buffer_t buffer, ggml_tensor * tensor, uint8_t value, size_t offset, size_t size) {
const size_t n_bufs = ggml_backend_meta_buffer_n_bufs(buffer);
const ggml_backend_meta_split_state split_state =
ggml_backend_meta_get_split_state(tensor, /*assume_sync =*/ false);
GGML_ASSERT(ggml_is_contiguous(tensor) || split_state.axis == GGML_BACKEND_SPLIT_AXIS_MIRRORED);
if (split_state.n_segments != 1 || split_state.nr[0] != 1) {
GGML_ASSERT(split_state.axis >= 0 && split_state.axis < GGML_MAX_DIMS);
GGML_ASSERT(split_state.nr[0] != 0);
GGML_ASSERT(tensor->ne[3] == 1);
std::vector<size_t> simple_offsets(n_bufs, 0);
if (split_state.axis == GGML_BACKEND_SPLIT_AXIS_0) {
GGML_ASSERT(tensor->ne[2] == 1);
const size_t row_stride = tensor->nb[1];
GGML_ASSERT(offset % row_stride == 0);
GGML_ASSERT(size % row_stride == 0);
const int64_t row_start = offset / row_stride;
const int64_t row_count = size / row_stride;
GGML_ASSERT(row_start + row_count <= tensor->ne[1]);
const int64_t blck_size = ggml_blck_size(tensor->type);
for (size_t s = 0; s < split_state.n_segments; s++) {
for (size_t r = 0; r < split_state.nr[s]; r++) {
for (size_t j = 0; j < n_bufs; j++) {
ggml_tensor * simple_tensor = ggml_backend_meta_buffer_simple_tensor(tensor, j);
GGML_ASSERT(split_state.ne[s*n_bufs + j] % blck_size == 0);
const size_t nbytes = split_state.ne[s*n_bufs + j]/blck_size * tensor->nb[0];
for (int64_t row = 0; row < row_count; row++) {
ggml_backend_tensor_memset(simple_tensor, value,
simple_offsets[j] + (row_start + row)*simple_tensor->nb[1], nbytes);
}
simple_offsets[j] += nbytes;
}
}
}
return;
}
GGML_ASSERT(split_state.axis == GGML_BACKEND_SPLIT_AXIS_1);
const size_t row_stride = tensor->nb[2];
GGML_ASSERT(offset % row_stride == 0);
GGML_ASSERT(size % row_stride == 0);
const int64_t row_start = offset / row_stride;
const int64_t row_count = size / row_stride;
GGML_ASSERT(row_start + row_count <= tensor->ne[2]);
for (size_t s = 0; s < split_state.n_segments; s++) {
for (size_t r = 0; r < split_state.nr[s]; r++) {
for (size_t j = 0; j < n_bufs; j++) {
ggml_tensor * simple_tensor = ggml_backend_meta_buffer_simple_tensor(tensor, j);
const size_t nbytes = split_state.ne[s*n_bufs + j] * tensor->nb[1];
for (int64_t row = 0; row < row_count; row++) {
ggml_backend_tensor_memset(simple_tensor, value,
simple_offsets[j] + (row_start + row)*simple_tensor->nb[2], nbytes);
}
simple_offsets[j] += nbytes;
}
}
}
return;
}
switch (split_state.axis) {
case GGML_BACKEND_SPLIT_AXIS_0:
case GGML_BACKEND_SPLIT_AXIS_1:
case GGML_BACKEND_SPLIT_AXIS_2: {
const size_t chunk_size_full = tensor->nb[split_state.axis + 1];
GGML_ASSERT(offset % chunk_size_full == 0);
GGML_ASSERT(size % chunk_size_full == 0);
const int64_t i_start = offset / chunk_size_full;
const int64_t i_stop = (offset + size) / chunk_size_full;
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 = simple_tensor->nb[split_state.axis + 1];
if (chunk_size == 0) {
continue;
}
for (int64_t i = i_start; i < i_stop; i++) {
ggml_backend_tensor_memset(simple_tensor, value, i*chunk_size, chunk_size);
}
}
} break;
case GGML_BACKEND_SPLIT_AXIS_PARTIAL: {
GGML_ASSERT(value == 0);
[[fallthrough]];
}
case GGML_BACKEND_SPLIT_AXIS_MIRRORED: {
for (size_t j = 0; j < n_bufs; j++) {
ggml_tensor * simple_tensor = ggml_backend_meta_buffer_simple_tensor(tensor, j);
ggml_backend_tensor_memset(simple_tensor, value, offset, size);
}
} break;
default: {
GGML_ABORT("fatal error");
}
}
}
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);
const ggml_backend_meta_split_state split_state = ggml_backend_meta_get_split_state(tensor, /*assume_sync =*/ false);
@@ -1352,15 +1504,29 @@ static void ggml_backend_meta_buffer_set_tensor(ggml_backend_buffer_t buffer, gg
} break;
case GGML_BACKEND_SPLIT_AXIS_PARTIAL: {
GGML_ASSERT(tensor->type == GGML_TYPE_F32);
const int64_t ne = ggml_nelements(tensor);
std::vector<float> tmp;
tmp.reserve(ne);
for (int64_t i = 0; i < ne; i++) {
tmp.push_back(((const float *) data)[i] / n_bufs);
GGML_ASSERT(offset % sizeof(float) == 0);
GGML_ASSERT(size % sizeof(float) == 0);
const size_t n_values = size / sizeof(float);
size_t n_contributors = 0;
for (size_t j = 0; j < n_bufs; j++) {
n_contributors += split_state.ne[j] != 0;
}
const bool has_contributor_mask = n_contributors != 0;
if (!has_contributor_mask) {
n_contributors = n_bufs;
}
std::vector<float> tmp(n_values);
for (size_t i = 0; i < n_values; i++) {
tmp[i] = ((const float *) data)[i] / n_contributors;
}
std::vector<float> zero;
if (has_contributor_mask) {
zero.resize(n_values, 0.0f);
}
for (size_t j = 0; j < n_bufs; j++) {
ggml_tensor * simple_tensor = ggml_backend_meta_buffer_simple_tensor(tensor, j);
ggml_backend_tensor_set(simple_tensor, tmp.data(), offset, size);
const float * partial = has_contributor_mask && split_state.ne[j] == 0 ? zero.data() : tmp.data();
ggml_backend_tensor_set(simple_tensor, partial, offset, size);
}
} break;
default: {
@@ -1488,7 +1654,7 @@ static const ggml_backend_buffer_i ggml_backend_meta_buffer_iface = {
/* .free_buffer = */ ggml_backend_meta_buffer_free_buffer,
/* .get_base = */ ggml_backend_meta_buffer_get_base,
/* .init_tensor = */ ggml_backend_meta_buffer_init_tensor,
/* .memset_tensor = */ nullptr, // TODO implement
/* .memset_tensor = */ ggml_backend_meta_buffer_memset_tensor,
/* .set_tensor = */ ggml_backend_meta_buffer_set_tensor,
/* .get_tensor = */ ggml_backend_meta_buffer_get_tensor,
/* .set_tensor_2d = */ nullptr,
@@ -1841,7 +2007,7 @@ static enum ggml_status ggml_backend_meta_graph_compute(ggml_backend_t backend,
{
// For MoE models it may make sense to delay the AllReduce in order to reduce I/O:
auto get_i_delayed = [&](const int i) -> int {
auto get_i_delayed_branch = [&](const int i) -> int {
int id = i; // i_delayed
int idr = i; // i_delayed return, last safe return value
@@ -1941,6 +2107,62 @@ static enum ggml_status ggml_backend_meta_graph_compute(ggml_backend_t backend,
return idr;
};
// AllReduce(a) + AllReduce(b) == AllReduce(a + b) for independent partial branches.
auto get_i_delayed = [&](const int i) -> int {
const int i_delayed = get_i_delayed_branch(i);
ggml_tensor * node = cgraph->nodes[i_delayed];
if (ggml_node_get_use_count(cgraph, i_delayed) != 1) {
return i_delayed;
}
for (int id = i_delayed + 1; id < cgraph->n_nodes; id++) {
ggml_tensor * next = cgraph->nodes[id];
if (next->view_src == node) {
return i_delayed;
}
for (int s = 0; s < GGML_MAX_SRC; s++) {
if (next->src[s] == node) {
return i_delayed;
}
}
if (next->view_src != nullptr && next->view_src->op == GGML_OP_NONE && ggml_backend_buffer_is_host(next->view_src->buffer)) {
continue;
}
if (ggml_backend_meta_get_split_state(next, false).axis != GGML_BACKEND_SPLIT_AXIS_PARTIAL) {
continue;
}
const int i_other = id;
const int i_other_delayed = get_i_delayed_branch(i_other);
ggml_tensor * other = cgraph->nodes[i_other_delayed];
if (ggml_node_get_use_count(cgraph, i_other_delayed) != 1 || i_other_delayed + 1 >= cgraph->n_nodes) {
return i_delayed;
}
ggml_tensor * sum = cgraph->nodes[i_other_delayed + 1];
if (sum->op != GGML_OP_ADD ||
!ggml_are_same_shape(node, other) || node->type != other->type || sum->type != node->type ||
!((sum->src[0] == node && sum->src[1] == other) ||
(sum->src[0] == other && sum->src[1] == node)) ||
ggml_backend_meta_get_split_state(sum, false).axis != GGML_BACKEND_SPLIT_AXIS_MIRRORED) {
return i_delayed;
}
for (size_t j = 0; j < n_backends; j++) {
auto & bcj = backend_ctx->backend_configs[j];
const bool compute = bcj.nodes[i]->flags & GGML_TENSOR_FLAG_COMPUTE;
const bool compute_other = bcj.nodes[i_other]->flags & GGML_TENSOR_FLAG_COMPUTE;
if (compute != compute_other) {
return i_delayed;
}
}
return i_other_delayed + 1;
}
return i_delayed;
};
int i_start = 0;
for (int i = 0; i < cgraph->n_nodes; i++) {
ggml_tensor * node = cgraph->nodes[i];
ggml/src/ggml-cuda/ggml-cuda.cu
+7 -2
View File
@@ -40,6 +40,7 @@ bool g_mul_mat_q = true;
#include "ggml-cuda/out-prod.cuh"
#include "ggml-cuda/pad.cuh"
#include "ggml-cuda/pool2d.cuh"
#include "ggml-cuda/pool1d.cuh"
#include "ggml-cuda/quantize.cuh"
#include "ggml-cuda/rope.cuh"
#include "ggml-cuda/roll.cuh"
@@ -2330,6 +2331,9 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg
case GGML_OP_POOL_2D:
ggml_cuda_op_pool2d(ctx, dst);
break;
case GGML_OP_POOL_1D:
ggml_cuda_op_pool1d(ctx, dst);
break;
case GGML_OP_SUM:
ggml_cuda_op_sum(ctx, dst);
break;
@@ -4620,8 +4624,8 @@ static std::string ggml_cuda_device_description(int device) {
const ggml_cuda_device_info & info = ggml_cuda_info();
std::string description = prop.name;
if (info.device_count > info.physical_device_count) {
description += " (physical device " + std::to_string(info.devices[device].physical_device) +
", virtual device " + std::to_string(info.devices[device].virtual_index) + ")";
description += " (dev p" + std::to_string(info.devices[device].physical_device) +
"/v" + std::to_string(info.devices[device].virtual_index) + ")";
}
return description;
}
@@ -5258,6 +5262,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
case GGML_OP_CONV_2D_DW:
return op->src[0]->type == GGML_TYPE_F32;
case GGML_OP_CONV_TRANSPOSE_2D:
case GGML_OP_POOL_1D:
case GGML_OP_POOL_2D:
return true;
case GGML_OP_ACC:
ggml/src/ggml-cuda/pool1d.cu
+85
View File
@@ -0,0 +1,85 @@
#include "pool1d.cuh"
static __global__ void pool1d_nchw_kernel(
const int iw, const int ow,
const int kw, const int sw, const int pw,
const int parallel_elements,
const float * src, float * dst, const enum ggml_op_pool op) {
const int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx >= parallel_elements) {
return;
}
const int nc = idx / ow;
const int cur_ow = idx % ow;
const float * i_ptr = src + nc * iw;
float * o_ptr = dst + nc * ow;
const int start = cur_ow * sw - pw;
const int b = max(0, start);
const int e = min(iw, start + kw);
float res;
switch (op) {
case GGML_OP_POOL_AVG: res = 0.0f; break;
case GGML_OP_POOL_MAX: res = -FLT_MAX; break;
default: return;
}
int count = 0;
for (int i = b; i < e; i++) {
#if __CUDA_ARCH__ >= 350
float cur = __ldg(i_ptr + i);
#else
float cur = i_ptr[i];
#endif
switch (op) {
case GGML_OP_POOL_AVG: res += cur; break;
case GGML_OP_POOL_MAX: res = max(res, cur); break;
default: break;
}
count++;
}
if (op == GGML_OP_POOL_AVG) {
res = (count > 0) ? (res / count) : 0.0f;
}
o_ptr[cur_ow] = res;
}
static void pool1d_nchw_kernel_f32_f32_cuda(
const int iw, const int ow,
const int kw, const int sw, const int pw,
const int parallel_elements,
const float * src, float * dst, const enum ggml_op_pool op,
cudaStream_t stream) {
const int num_blocks = (parallel_elements + CUDA_POOL1D_BLOCK_SIZE - 1) / CUDA_POOL1D_BLOCK_SIZE;
dim3 block_nums(num_blocks);
pool1d_nchw_kernel<<<block_nums, CUDA_POOL1D_BLOCK_SIZE, 0, stream>>>(iw, ow, kw, sw, pw, parallel_elements, src, dst, op);
}
void ggml_cuda_op_pool1d(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
const ggml_tensor * src0 = dst->src[0];
const float * src0_d = (const float *)src0->data;
float * dst_d = (float *)dst->data;
cudaStream_t stream = ctx.stream();
GGML_ASSERT(src0->type == GGML_TYPE_F32);
GGML_ASSERT( dst->type == GGML_TYPE_F32);
const int32_t * opts = (const int32_t *)dst->op_params;
enum ggml_op_pool op = static_cast<ggml_op_pool>(opts[0]);
const int k0 = opts[1];
const int s0 = opts[2];
const int p0 = opts[3];
const int64_t IW = src0->ne[0];
const int64_t OW = dst->ne[0];
const int64_t nr = ggml_nrows(src0);
const int parallel_elements = (int)(nr * OW);
pool1d_nchw_kernel_f32_f32_cuda(IW, OW, k0, s0, p0, parallel_elements, src0_d, dst_d, op, stream);
}
ggml/src/ggml-cuda/pool1d.cuh
+5
View File
@@ -0,0 +1,5 @@
#include "common.cuh"
#define CUDA_POOL1D_BLOCK_SIZE 256
void ggml_cuda_op_pool1d(ggml_backend_cuda_context & ctx, ggml_tensor * dst);
ggml/src/ggml-metal/ggml-metal-device.cpp
+2 -2
View File
@@ -17,10 +17,10 @@ struct ggml_metal_device_deleter {
typedef std::unique_ptr<ggml_metal_device, ggml_metal_device_deleter> ggml_metal_device_ptr;
ggml_metal_device_t ggml_metal_device_get(int device) {
ggml_metal_device_t ggml_metal_device_get(int device, int n_devices) {
static std::vector<ggml_metal_device_ptr> devs;
devs.emplace_back(ggml_metal_device_init(device));
devs.emplace_back(ggml_metal_device_init(device, n_devices));
return devs.back().get();
}
ggml/src/ggml-metal/ggml-metal-device.h
+4 -2
View File
@@ -259,6 +259,8 @@ enum ggml_metal_device_id {
struct ggml_metal_device_props {
int device;
int device_phys;
int device_virt;
char name[128];
char desc[128];
@@ -286,10 +288,10 @@ typedef struct ggml_metal_event * ggml_metal_event_t;
void ggml_metal_event_encode_signal(ggml_metal_event_t ev, ggml_metal_cmd_buf_t cmd_buf);
void ggml_metal_event_encode_wait (ggml_metal_event_t ev, ggml_metal_cmd_buf_t cmd_buf);
ggml_metal_device_t ggml_metal_device_init(int device);
ggml_metal_device_t ggml_metal_device_init(int device, int n_devices);
void ggml_metal_device_free(ggml_metal_device_t dev);
ggml_metal_device_t ggml_metal_device_get(int device);
ggml_metal_device_t ggml_metal_device_get(int device, int n_devices);
void * ggml_metal_device_get_obj (ggml_metal_device_t dev); // id<MTLDevice>
void * ggml_metal_device_get_queue(ggml_metal_device_t dev); // id<MTLCommandQueue>
ggml/src/ggml-metal/ggml-metal-device.m
+14 -2
View File
@@ -717,7 +717,7 @@ static enum ggml_metal_device_id ggml_metal_device_id_parse(const char * name) {
return GGML_METAL_DEVICE_GENERIC;
}
ggml_metal_device_t ggml_metal_device_init(int device) {
ggml_metal_device_t ggml_metal_device_init(int device, int n_devices) {
ggml_metal_device_t dev = calloc(1, sizeof(struct ggml_metal_device));
assert(dev != NULL);
@@ -734,6 +734,12 @@ ggml_metal_device_t ggml_metal_device_init(int device) {
dev->addr_virt = 0x000000400ULL;
dev->props.device = device;
// the Metal backend uses the system default device as the single physical device;
// additional (virtual) devices are emulated on top of it via GGML_METAL_DEVICES
dev->props.device_phys = 0;
dev->props.device_virt = device;
dev->props.has_simdgroup_reduction = [dev->mtl_device supportsFamily:MTLGPUFamilyApple7];
dev->props.has_simdgroup_reduction |= [dev->mtl_device supportsFamily:MTLGPUFamilyMetal3_GGML];
@@ -897,7 +903,13 @@ ggml_metal_device_t ggml_metal_device_init(int device) {
}
snprintf(dev->props.name, sizeof(dev->props.name), "%s%d", "MTL", device);
snprintf(dev->props.desc, sizeof(dev->props.desc), "%s", [[dev->mtl_device name] UTF8String]);
const char * gpu_name = [[dev->mtl_device name] UTF8String];
if (n_devices > 1) {
snprintf(dev->props.desc, sizeof(dev->props.desc), "%s (dev p%d/v%d)",
gpu_name, dev->props.device_phys, dev->props.device_virt);
} else {
snprintf(dev->props.desc, sizeof(dev->props.desc), "%s", gpu_name);
}
dev->library = ggml_metal_library_init(dev);
if (!dev->library) {
ggml/src/ggml-metal/ggml-metal.cpp
+1 -1
View File
@@ -891,7 +891,7 @@ static ggml_backend_dev_t ggml_backend_metal_device_init(ggml_backend_reg_t reg,
return new ggml_backend_device {
/* .iface = */ ggml_backend_metal_device_i,
/* .reg = */ reg,
/* .context = */ ggml_metal_device_get(device),
/* .context = */ ggml_metal_device_get(device, g_devices),
};
}
ggml/src/ggml-vulkan/ggml-vulkan.cpp
+26
View File
@@ -961,6 +961,7 @@ struct vk_device_struct {
vk_pipeline pipeline_diag[2];
vk_pipeline pipeline_clamp[2];
vk_pipeline pipeline_pad_f32;
vk_pipeline pipeline_pad_reflect_1d_f32;
vk_pipeline pipeline_roll_f32;
vk_pipeline pipeline_repeat_i32, pipeline_repeat_back_f32;
vk_pipeline pipeline_repeat_i16;
@@ -5636,6 +5637,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
ggml_vk_create_pipeline(device, device->pipeline_diag[1], "diag_f16", diag_f16_len, diag_f16_data, "main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1);
ggml_vk_create_pipeline(device, device->pipeline_pad_f32, "pad_f32", pad_f32_len, pad_f32_data, "main", 2, sizeof(vk_op_pad_push_constants), {512, 1, 1}, {}, 1);
ggml_vk_create_pipeline(device, device->pipeline_pad_reflect_1d_f32, "pad_reflect_1d_f32", pad_reflect_1d_f32_len, pad_reflect_1d_f32_data, "main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1);
ggml_vk_create_pipeline(device, device->pipeline_roll_f32, "roll_f32", roll_f32_len, roll_f32_data, "main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1);
@@ -11369,6 +11371,11 @@ static vk_pipeline ggml_vk_op_get_pipeline(ggml_backend_vk_context * ctx, const
return ctx->device->pipeline_pad_f32;
}
return nullptr;
case GGML_OP_PAD_REFLECT_1D:
if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
return ctx->device->pipeline_pad_reflect_1d_f32;
}
return nullptr;
case GGML_OP_ROLL:
if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
return ctx->device->pipeline_roll_f32;
@@ -12272,6 +12279,7 @@ static void ggml_vk_op_f32(ggml_backend_vk_context * ctx, vk_context& subctx, co
case GGML_OP_CLAMP:
case GGML_OP_LEAKY_RELU:
case GGML_OP_PAD:
case GGML_OP_PAD_REFLECT_1D:
case GGML_OP_ROLL:
case GGML_OP_REPEAT:
case GGML_OP_REPEAT_BACK:
@@ -13144,6 +13152,17 @@ static void ggml_vk_pad(ggml_backend_vk_context * ctx, vk_context& subctx, const
ggml_vk_op_f32(ctx, subctx, src0, nullptr, nullptr, nullptr, dst, GGML_OP_PAD, std::move(p));
}
static void ggml_vk_pad_reflect_1d(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst) {
const uint32_t p0 = (uint32_t)dst->op_params[0];
const uint32_t p1 = (uint32_t)dst->op_params[1];
vk_op_unary_push_constants p = vk_op_unary_push_constants_init(src0, dst, ggml_nelements(dst));
memcpy(&p.param1, &p0, sizeof(float));
memcpy(&p.param2, &p1, sizeof(float));
ggml_vk_op_f32(ctx, subctx, src0, nullptr, nullptr, nullptr, dst, GGML_OP_PAD_REFLECT_1D, std::move(p));
}
static void ggml_vk_roll(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst) {
const int32_t s0 = ggml_get_op_params_i32(dst, 0);
const int32_t s1 = ggml_get_op_params_i32(dst, 1);
@@ -15553,6 +15572,10 @@ static bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgr
case GGML_OP_PAD:
ggml_vk_pad(ctx, compute_ctx, src0, node);
break;
case GGML_OP_PAD_REFLECT_1D:
ggml_vk_pad_reflect_1d(ctx, compute_ctx, src0, node);
break;
case GGML_OP_ROLL:
ggml_vk_roll(ctx, compute_ctx, src0, node);
@@ -18479,6 +18502,7 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
case GGML_OP_SCALE:
return ggml_is_contiguous(op->src[0]) && op->src[0]->type == GGML_TYPE_F32;
case GGML_OP_PAD:
case GGML_OP_PAD_REFLECT_1D:
case GGML_OP_ROLL:
return op->src[0]->type == GGML_TYPE_F32;
case GGML_OP_DIAG_MASK_INF:
@@ -19261,6 +19285,8 @@ static void ggml_vk_check_results_0(ggml_backend_vk_context * ctx, ggml_cgraph *
} else if (tensor->op == GGML_OP_PAD) {
tensor_clone = ggml_pad_ext(ggml_ctx, src_clone[0], tensor->op_params[0], tensor->op_params[1], tensor->op_params[2], tensor->op_params[3],
tensor->op_params[4], tensor->op_params[5], tensor->op_params[6], tensor->op_params[7]);
} else if (tensor->op == GGML_OP_PAD_REFLECT_1D) {
tensor_clone = ggml_pad_reflect_1d(ggml_ctx, src_clone[0], tensor->op_params[0], tensor->op_params[1]);
} else if (tensor->op == GGML_OP_REPEAT) {
tensor_clone = ggml_repeat(ggml_ctx, src_clone[0], tensor);
} else if (tensor->op == GGML_OP_REPEAT_BACK) {
ggml/src/ggml-vulkan/vulkan-shaders/pad_reflect_1d.comp
+43
View File
@@ -0,0 +1,43 @@
#version 450
#include "types.glsl"
#include "generic_unary_head.glsl" // included to use functions like fastdiv etc.
layout(local_size_x = 512, local_size_y = 1, local_size_z = 1) in;
void main() {
const uint idx = get_idx();
if (idx >= p.ne) {
return;
}
const uint p0 = floatBitsToUint(p.param1);
const uint p1 = floatBitsToUint(p.param2);
const uint i3 = fastdiv(idx, p.ne1_012mp, fastdiv_L(p.ne1_Ls, 0));
const uint i3_offset = i3 * p.ne12 * p.ne11 * p.ne10;
const uint i2 = fastdiv(idx - i3_offset, p.ne1_01mp, fastdiv_L(p.ne1_Ls, 1));
const uint i2_offset = i2 * p.ne11 * p.ne10;
const uint i1 = fastdiv(idx - i3_offset - i2_offset, p.ne1_0mp, fastdiv_L(p.ne1_Ls, 2));
const uint i0 = idx - i3_offset - i2_offset - i1 * p.ne10;
uint src_col;
if (i0 < p0) {
src_col = p0 - i0; // left pad area
} else if (i0 < p0 + p.ne00) {
src_col = i0 - p0; // center area
} else {
src_col = 2u * p.ne00 - 2u - (i0 - p0); // right pad area
}
const uint src_idx = i3 * p.nb03 + i2 * p.nb02 + i1 * p.nb01 + src_col * p.nb00;
const uint d_idx = i3 * p.nb13 + i2 * p.nb12 + i1 * p.nb11 + i0 * p.nb10;
// copy the computed value to the destination tensor
data_d[get_doffset() + d_idx] = D_TYPE(data_a[get_aoffset() + src_idx]);
}
ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp
+1
View File
@@ -922,6 +922,7 @@ void process_shaders() {
string_to_spv("scale_f32", "scale.comp", {{"A_TYPE", "float"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}});
string_to_spv("pad_f32", "pad.comp", {{"A_TYPE", "float"}, {"D_TYPE", "float"}});
string_to_spv("pad_reflect_1d_f32", "pad_reflect_1d.comp", {{"A_TYPE", "float"}, {"D_TYPE", "float"}});
string_to_spv("concat_i8", "concat.comp", {{"A_TYPE", "uint8_t"}, {"B_TYPE", "uint8_t"}, {"D_TYPE", "uint8_t"}});
string_to_spv("concat_i16", "concat.comp", {{"A_TYPE", "uint16_t"}, {"B_TYPE", "uint16_t"}, {"D_TYPE", "uint16_t"}});
ggml/src/ggml.c
+35 -19
View File
@@ -4058,6 +4058,41 @@ struct ggml_tensor * ggml_diag_mask_zero_inplace(
return ggml_diag_mask_zero_impl(ctx, a, n_past, true);
}
// ggml_clamp
static struct ggml_tensor * ggml_clamp_impl(
struct ggml_context * ctx,
struct ggml_tensor * a,
float min,
float max,
bool inplace) {
struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
float params[] = { min, max };
ggml_set_op_params(result, params, sizeof(params));
result->op = GGML_OP_CLAMP;
result->src[0] = a;
return result;
}
struct ggml_tensor * ggml_clamp(
struct ggml_context * ctx,
struct ggml_tensor * a,
float min,
float max) {
return ggml_clamp_impl(ctx, a, min, max, false);
}
struct ggml_tensor * ggml_clamp_inplace(
struct ggml_context * ctx,
struct ggml_tensor * a,
float min,
float max) {
return ggml_clamp_impl(ctx, a, min, max, true);
}
// ggml_soft_max
static struct ggml_tensor * ggml_soft_max_impl(
@@ -4454,25 +4489,6 @@ struct ggml_tensor * ggml_rope_set_offset(
return a;
}
// ggml_clamp
struct ggml_tensor * ggml_clamp(
struct ggml_context * ctx,
struct ggml_tensor * a,
float min,
float max) {
// TODO: when implement backward, fix this:
struct ggml_tensor * result = ggml_view_tensor(ctx, a);
float params[] = { min, max };
ggml_set_op_params(result, params, sizeof(params));
result->op = GGML_OP_CLAMP;
result->src[0] = a;
return result;
}
static int64_t ggml_calc_conv_output_size(int64_t ins, int64_t ks, int s, int p, int d) {
return (ins + 2 * p - d * (ks - 1) - 1) / s + 1;
}
gguf-py/gguf/constants.py
+1 -1
View File
@@ -3822,7 +3822,7 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
MODEL_TENSOR.FFN_DOWN_SHEXP,
MODEL_TENSOR.FFN_UP_SHEXP,
MODEL_TENSOR.FFN_EXP_PROBS_B,
# NextN/MTP tensors - preserved but unused
# NextN/MTP tensors
MODEL_TENSOR.NEXTN_EH_PROJ,
MODEL_TENSOR.NEXTN_EMBED_TOKENS,
MODEL_TENSOR.NEXTN_ENORM,
gpttype_adapter.cpp
+3 -3
View File
@@ -5107,7 +5107,7 @@ std::string gpttype_parse_chat_tool_calls(const std::string & generated_text,
return "";
}
json tools = json::parse(tools_json);
common_json tools = common_json::parse(tools_json);
if(!tools.is_array() || tools.empty())
{
return "";
@@ -5135,7 +5135,7 @@ std::string gpttype_parse_chat_tool_calls(const std::string & generated_text,
if(!chat_template_kwargs_json.empty())
{
json kwargs = json::parse(chat_template_kwargs_json);
common_json kwargs = common_json::parse(chat_template_kwargs_json);
if(kwargs.is_object())
{
for(const auto & item : kwargs.items())
@@ -5181,7 +5181,7 @@ std::string gpttype_parse_chat_tool_calls(const std::string & generated_text,
return "";
}
json tool_calls = parsed.to_json_oaicompat().value("tool_calls", json::array());
common_json tool_calls = parsed.to_json_oaicompat().value("tool_calls", common_json::array());
return tool_calls.dump();
}
catch(const std::exception & e)
include/llama.h
+1 -1
View File
@@ -736,7 +736,7 @@ extern "C" {
// Removes all tokens that belong to the specified sequence and have positions in [p0, p1)
// Returns false if a partial sequence cannot be removed. Removing a whole sequence never fails
// seq_id < 0 : match any sequence
// seq_id < 0 : match any sequence [TAG_LLAMA_SEQ_ID_NEG]
// p0 < 0 : [0, p1]
// p1 < 0 : [p0, inf)
LLAMA_API bool llama_memory_seq_rm(
skills/add-new-model/SKILL.md
-99
View File
@@ -1,99 +0,0 @@
---
name: add-new-model
description: Guided workflow for adding a new model architecture to llama.cpp. Use when the user wants to add/port a new model architecture.
---
# Add a new model architecture to llama.cpp
This skill walks a contributor through adding a new model architecture. AI-generated code is permitted in this project, so you may write full implementations for the steps below rather than only pointing at patterns - but follow `AGENTS.md`'s AI usage policy throughout:
- The contributor is 100% responsible for every line, however it was produced. They must be able to explain and defend any part of it to a reviewer. Check in with them as you go (don't silently generate everything and hand over a finished diff) so they actually absorb what was written.
- Before writing code, make sure the contributor owns the design choices for this architecture (which reference model to follow, how non-standard bits like RoPE variants or MoE routing should be handled) - AI accelerates a design the contributor has already made, it doesn't make the design for them.
- Disclosure is mandatory: any AI-meaningful contribution must be disclosed per the PR template. Remind the contributor of this before they open the PR.
- Never write the PR description, commit message, GitHub issue/discussion post, or reviewer replies - those must come from the contributor. If asked to commit on their behalf, use `Assisted-by:` (never `Co-authored-by:`) and only after explicit confirmation.
- If the requested change looks large or introduces a new pattern not covered here, pause and tell the user this kind of change is likely to need prior discussion with maintainers before a PR.
- Keep the PR self-contained. If the work would require a lot of unconventional changes outside the new model file(s) (e.g. touching shared graph-building code, the sampler, or core APIs in ways other models don't), STOP and tell the contributor to open a discussion/issue first - invasive or excessive changes get closed without full review.
- Do not bundle unrelated work into this PR - see Step 4 and Step 5 below for the specifics on multimodal and chat-template/parsing work.
- Never hack around RoPE with a custom sin/cos implementation. Several past PRs tried this and were closed. If the existing `ggml_rope_ext` (see Step 2's RoPE tips) genuinely cannot express what this model needs, the contributor should open an issue to discuss it with maintainers first - not send a PR with a custom RoPE implementation.
Before starting, read `CONTRIBUTING.md`, `AGENTS.md` and `docs/development/HOWTO-add-model.md` if they are not already in context. Also run `git log --oneline -- src/models` and look at at least 3 recent PRs that added a model (their merge commits/diffs) - this shows current convention more reliably than the docs, which can lag behind.
## Step 0 - Scope and dedup check
Ask the contributor:
1. Which model (HF repo id or name)? Is it text-only or does it have a multimodal (vision/audio) encoder?
2. Do they already have the HF `config.json`/weights available locally?
3. Have they checked for an existing PR/issue on this model? Suggest `gh search issues "<model name>"` and `gh search prs "<model name>"` in the `ggml-org/llama.cpp` repo. If an existing PR covers it, the contributor should comment there and collaborate rather than open a duplicate (per CONTRIBUTING.md's AI Usage Policy).
4. What existing supported architecture is this model closest to (e.g. "Llama-like with sliding window", "MoE like DBRX", "BERT-style encoder")?
If the contributor doesn't know the closest reference architecture, you may grep `conversion/*.py` and `src/models/*.cpp` for architectures with a similar config shape (layer count, head count, MoE expert count, norm placement) and suggest 1-2 candidates - but let the contributor confirm the choice rather than picking one yourself; this choice is a design decision they need to own.
Do not proceed to Step 1 until the contributor has answered these and named a reference architecture.
## Step 1 - Convert the model to GGUF
Follow HOWTO-add-model.md section 1 for the actual touch points (conversion class registration, `constants.py`, `tensor_mapping.py`, etc.) - don't re-derive them here, read them from the doc.
Skill-specific addition: for each touch point, show the contributor the equivalent code in the reference architecture they named in Step 0 before writing the new version, and check that they understand what's different about their model (e.g. non-standard tensor shapes, extra hparams) rather than just copying the pattern silently.
## Step 2 - Define the architecture in llama.cpp
Follow HOWTO-add-model.md section 2 for the actual touch points (`llm_arch` enum, `LLM_ARCH_NAMES`, hparam loading, RoPE type case, etc.), including its "Tips and tricks" section for `ggml_rope_ext` gotchas.
Skill-specific addition: never hack around RoPE with a custom sin/cos implementation - see the RoPE rule above.
## Step 3 - Build the GGML graph
Follow HOWTO-add-model.md section 3 for the actual touch points (`src/models/<name>.cpp` struct, `llama_model_mapping` registration, etc.).
Skill-specific addition: before writing `src/models/<name>.cpp`, read at least 10 other files under `src/models/` (pick a mix, not just the one reference architecture) to confirm the struct layout, naming, and style you're about to write actually matches current convention - the pattern drifts over time and the HOWTO doc can lag behind it.
## Step 4 - Optional: multimodal encoder
Only do this if the contributor flagged a vision/audio encoder in Step 0. Follow HOWTO-add-model.md section 4 and `docs/multimodal.md` for the actual touch points (`MmprojModel` subclass, `clip.cpp`, `mtmd.cpp`, encoder graph in `tools/mtmd/models`, etc.).
Skill-specific addition, and read this carefully: **whether the multimodal encoder can be bundled into the same PR as the base text-model support depends on how conventional the change is.** It's OK to bundle it if the encoder support is conventional - i.e. no new infra or logic is needed, it's just a new cgraph reusing existing preprocessing/projector machinery (e.g. siglip/pixtral/qwen with just a new projector). If it requires anything beyond that - a new preprocessor, non-standard projector logic, or changes to shared `libmtmd` infra/logic - STOP, tell the contributor this is non-conventional, and have them land the text model first with the encoder as a dedicated follow-up PR. Do not let this decision pass silently - call it out explicitly to the contributor before writing any `clip.cpp`/`mtmd.cpp` code.
## Step 5 - Optional: chat template / parsing support
Only do this if the model needs a new built-in chat template (`src/llama-chat.cpp`) or a new output parser (see `docs/development/parsing.md` and `docs/autoparser.md`). If either is needed beyond what a user-supplied Jinja template already covers, treat it as its own dedicated follow-up PR, not part of the base model-support PR - call this out explicitly to the contributor rather than silently bundling it in.
## Common pitfalls (from past PR reviews)
These recur often enough in review comments on past add-model PRs that they're worth checking proactively, not just waiting for a reviewer to catch them:
- Don't validate the same hparam/config assumption in both the Python conversion script and the C++ load path - pick one layer to own the check, duplicating it just adds maintenance surface.
- Optional hparams that are genuinely absent from some configs (e.g. a shared-expert count) should be read with an explicit optional/fallback accessor, not assumed present.
- Hparams that are actually load-bearing (the model produces wrong output or crashes without them, e.g. `sliding_window_pattern`, norm-eps) must hard-error if missing, not silently fall back to a default.
- Don't bake a default chat template into the C++ binary - inject it into the GGUF at conversion time instead, since one `llm_arch` can be reused by multiple fine-tunes with different templates, and a baked-in C++ default fails silently for those.
- Before writing a dedicated tool-call/output parser, check whether the existing autoparser already handles the template (`llama-debug-template-parser <jinja>` shows what it detects).
- Marking a custom EOS/closing-tag token as `eot` at conversion time isn't always sufficient - in long/agentic generations a model can emit the closing sequence as literal text instead of the token, so generation never stops on EOG and raw text leaks past the parser. Verify this case, not just the token path.
- If reusing or aliasing an existing pre-tokenizer for convenience, justify and test that choice explicitly - silent reuse is an easy source of subtle tokenizer bugs.
- Watch for excessive graph splits caused by building per-layer view/index tensors inside the layer loop - hoist tensors that don't vary per layer out of the loop (relevant if you hit `GGML_SCHED_MAX_SPLIT_INPUTS`).
- A custom KQ mask fed into flash attention must match FA's expected dtype - cast it to F16 before passing it to `build_attn_mha` when FA is enabled.
- When padding a custom KV-cache size to an alignment (e.g. `GGML_PAD(..., 256)`), apply the padding after all other size adjustments, not before - otherwise later logic can un-align it again.
- For non-standard cache/SWA (sliding-window-attention) semantics, override the dedicated hook (e.g. `llama_model_n_swa()`) rather than mutating hparams to fake the behavior - hparams may be read elsewhere for unrelated purposes.
- Don't ship unfinished or unverified speculative-decoding (e.g. MTP) scaffolding in the base model PR - if it hasn't actually been confirmed to work, pull it out and land it as its own follow-up.
- Conversion code should call into the base class's existing hparam logic (e.g. `super().set_gguf_parameters()`) rather than re-deriving it - large blocks of code that duplicate what `TextModel`/`MmprojModel` already provide will get flagged as redundant.
- Do constant tensor modifications (e.g. `norm(1 + weight)`) and permutations/chunking at conversion time, not in the graph - see HOWTO-add-model.md's "Prefer conversion-time tensor modifications" tip (Gemma 3 folds its `1 +` into the weights, Qwen3-Next permutes in `modify_tensors`). Doing these at runtime in the graph is very likely to be rejected as over-complicated; if you genuinely can't do it at conversion time, open a discussion first explaining why rather than implementing it in the graph.
- Exception: a plain `weight * scale` with a constant scale is usually better applied at inference time instead of being folded into the weight at conversion. The scale conceptually applies to the activation, not the weight, so folding it in can hurt numerical stability, and it shifts the weight's value range in a way that can make quantization worse.
## Validation checklist
Reference: `examples/model-conversion/README.md`.
1. Convert to GGUF, then inspect/run both the original and converted tensors.
2. Run logits verification (original vs converted). If this model is a new version of an already-supported family, verify the *previous* version still passes logits verification first - numerical differences may be pre-existing, not caused by the new work. The tools to perform full logits validation are available in `examples/model-conversion`.
3. Quantize (including QAT variants if relevant) and re-verify.
4. Run perplexity evaluation (simple and full).
5. Sanity-check across `tools/cli`, `tools/completion`, `tools/imatrix`, `tools/quantize`, and `tools/server`.
6. CPU backend first; other backends (CUDA, Metal, ...) can be separate follow-up PRs per `CONTRIBUTING.md`.
7. Re-review every changed file against the coding/naming guidelines in `AGENTS.md` (and `CONTRIBUTING.md`'s "Coding guidelines"/"Naming guidelines" sections) - this is a separate pass from functional testing and is just as important: no forced line-wrapping, no unicode punctuation, minimal/non-redundant comments, `snake_case` naming (`kebab-case` for file names), matching indentation/brace style, etc.
## Before opening a PR
- Run the `code-review` skill on the diff first - it catches the convention and scope issues reviewers flag most often, and it's recommended to do this locally before pushing the PR.
- Confirm the contributor can explain every changed line to a reviewer and is prepared to be asked about any of it - this is required regardless of how much of the code was AI-generated.
- Confirm they did a comprehensive manual review of the full diff, not just a skim.
- Fill in the AI-disclosure section of `.github/pull_request_template.md` describing how AI was used (do not omit or understate this).
- Do not write the PR description, commit message, GitHub issue/discussion text, or any reviewer replies yourself - the contributor writes these.
src/llama-arch.cpp
-1
View File
@@ -1060,7 +1060,6 @@ bool llm_arch_supports_sm_tensor(const llm_arch & arch) {
case LLM_ARCH_OLMOE:
case LLM_ARCH_DEEPSEEK2:
case LLM_ARCH_DEEPSEEK32:
case LLM_ARCH_DEEPSEEK4:
case LLM_ARCH_DOTS3NOTE:
case LLM_ARCH_GLM_DSA:
case LLM_ARCH_BITNET:
src/llama-context.cpp
-4
View File
@@ -3228,8 +3228,6 @@ size_t llama_context::state_read_data(llama_io_read_i & io) {
}
size_t llama_context::state_seq_write_data(llama_io_write_i & io, llama_seq_id seq_id, llama_state_seq_flags flags) {
GGML_UNUSED(seq_id);
if (memory) {
memory->state_write(io, seq_id, flags);
}
@@ -3238,8 +3236,6 @@ size_t llama_context::state_seq_write_data(llama_io_write_i & io, llama_seq_id s
}
size_t llama_context::state_seq_read_data(llama_io_read_i & io, llama_seq_id seq_id, llama_state_seq_flags flags) {
GGML_UNUSED(seq_id);
if (memory) {
memory->state_read(io, seq_id, flags);
}
src/llama-kv-cache-dsv4.cpp
+98 -40
View File
@@ -599,6 +599,33 @@ static llama_kv_cache_dsv4_context::comp_plan dsv4_build_comp_plan(
}
}
if (ratio == DSV4_HCA_RATIO && !plan.state_pos.empty() && plan.state_write_idxs.empty()) {
assert(kv_size > 0);
// the last slot must not be live, or the dummy write would corrupt it;
// a full stream implies a completed block, which implies real writes
assert(plan.n_kv < (int64_t) kv_size);
// Keep the compress/write ops in the graph when no HCA block completes
// in this ubatch. The dummy block writes to the last cache slot and is
// masked out.
uint32_t i = 0;
while (i < ubatch.n_tokens && ubatch.pos[i] < 0) {
++i;
}
assert(i < ubatch.n_tokens);
const llama_seq_id seq_id = ubatch.seq_id[i][0];
const int64_t cache_off = dsv4_stream_offset(n_stream, seq_id, kv_size);
const int32_t source_idx = state_source_idx(seq_id, ubatch.pos[i]);
plan.state_write_idxs.push_back(cache_off + kv_size - 1);
plan.state_write_pos .push_back(0);
for (uint32_t j = 0; j < ratio; ++j) {
plan.state_read_idxs.push_back(source_idx);
}
}
if (overlap) {
// [ all blocks' prev-window indices | all blocks' cur-window indices ]
plan.state_read_idxs.reserve(overlap_prev_reads.size() + overlap_cur_reads.size());
@@ -608,7 +635,10 @@ static llama_kv_cache_dsv4_context::comp_plan dsv4_build_comp_plan(
overlap_cur_reads.begin(), overlap_cur_reads.end());
}
plan.n_kv = GGML_PAD(plan.n_kv, 256u);
// Keep the mask (and with it the compressed-attention branch) present even
// before the first block is visible, so the graph topology never changes.
// Padded slots are masked out; comp cache buffers are zero-initialized.
plan.n_kv = std::max<int64_t>(GGML_PAD(plan.n_kv, 256u), 256);
std::sort(persist_rows.begin(), persist_rows.end(),
[](const persist_row & a, const persist_row & b) {
@@ -620,16 +650,26 @@ static llama_kv_cache_dsv4_context::comp_plan dsv4_build_comp_plan(
plan.state_persist_dst_idxs.push_back(row.dst);
}
if (n_rs_seq > 0) {
for (uint32_t s = 0; s < ubatch.n_seqs_unq; ++s) {
const llama_seq_id seq_id = ubatch.seq_id_unq[s];
if (seq_id < 0 || (uint32_t) seq_id >= n_stream) {
continue;
// Emit restore/snapshot entries for all layout streams so that the
// graph tensor sizes do not depend on the ubatch's sequence count.
// Streams not present in the ubatch get no-op entries.
for (uint32_t stream = 0; stream < n_stream; ++stream) {
llama_seq_id seq_id = -1;
if (n_stream == 1) {
// a unified stream serves any single sequence
seq_id = ubatch.n_seqs_unq > 0 ? ubatch.seq_id_unq[0] : -1;
} else {
for (uint32_t s = 0; s < ubatch.n_seqs_unq; ++s) {
if (ubatch.seq_id_unq[s] == (llama_seq_id) stream) {
seq_id = ubatch.seq_id_unq[s];
break;
}
}
}
const int64_t stream_off = dsv4_stream_offset(n_stream, seq_id, state_size);
const uint32_t rollback = (uint32_t) seq_id < rs_idx.size() ? rs_idx[seq_id] : 0;
const int64_t stream_off = (int64_t) stream*state_size;
const uint32_t rollback = seq_id >= 0 && (uint32_t) seq_id < rs_idx.size() ? rs_idx[seq_id] : 0;
// Keep the restore graph fixed-width when no rollback is pending.
const int64_t src_plane = rollback > 0 && rollback <= n_rs_seq ? (int64_t) rollback*state_rows : 0;
for (uint32_t r = 0; r < state_size; ++r) {
@@ -639,35 +679,33 @@ static llama_kv_cache_dsv4_context::comp_plan dsv4_build_comp_plan(
std::vector<uint32_t> token_idxs;
token_idxs.reserve(ubatch.n_tokens);
for (uint32_t i = 0; i < ubatch.n_tokens; ++i) {
if (dsv4_token_has_seq(ubatch, i, seq_id)) {
token_idxs.push_back(i);
if (seq_id >= 0) {
for (uint32_t i = 0; i < ubatch.n_tokens; ++i) {
if (dsv4_token_has_seq(ubatch, i, seq_id)) {
token_idxs.push_back(i);
}
}
}
if (token_idxs.empty()) {
continue;
}
const uint32_t n_seq_tokens = (uint32_t) token_idxs.size();
const int64_t scratch_off = (int64_t) state_rows*(1 + n_rs_seq);
for (uint32_t d = 1; d <= n_rs_seq; ++d) {
const int64_t dst_plane = (int64_t) d*state_rows;
const uint32_t prefix = d <= n_seq_tokens ? n_seq_tokens - d : 0;
for (uint32_t r = 0; r < state_size; ++r) {
int32_t src;
if (d <= n_seq_tokens) {
const uint32_t prefix = n_seq_tokens - d;
src = (int32_t) (stream_off + r);
int32_t src = (int32_t) (stream_off + r);
for (uint32_t j = 0; j < prefix; ++j) {
const uint32_t i_tok = token_idxs[j];
if (ubatch.pos[i_tok] >= 0 && (uint32_t) (ubatch.pos[i_tok]%state_size) == r) {
src = (int32_t) (scratch_off + i_tok);
}
for (uint32_t j = 0; j < prefix; ++j) {
const uint32_t i_tok = token_idxs[j];
if (ubatch.pos[i_tok] >= 0 && (uint32_t) (ubatch.pos[i_tok]%state_size) == r) {
src = (int32_t) (scratch_off + i_tok);
}
} else {
const int64_t src_plane = (int64_t) (d - n_seq_tokens)*state_rows;
src = (int32_t) (src_plane + stream_off + r);
}
if (n_seq_tokens == 0) {
// no-op: copy the snapshot plane onto itself
src = (int32_t) (dst_plane + stream_off + r);
}
plan.state_snapshot_src_idxs.push_back(src);
@@ -683,10 +721,16 @@ static llama_kv_cache_dsv4_context::comp_plan dsv4_build_comp_plan(
}();
if (debug) {
LLAMA_LOG_INFO("%s: ratio=%u, n_tokens=%u, state_persist_dst=%s, state_write_pos=%s\n",
__func__, ratio, ubatch.n_tokens,
LLAMA_LOG_DEBUG("%s: ratio=%u, n_tokens=%u, n_seqs_unq=%u, state_persist_dst=%s, state_write_pos=%s\n",
__func__, ratio, ubatch.n_tokens, ubatch.n_seqs_unq,
dsv4_plan_positions(plan.state_persist_dst_idxs).c_str(),
dsv4_plan_positions(plan.state_write_pos).c_str());
for (uint32_t s = 0; s < ubatch.n_seqs_unq; ++s) {
const llama_seq_id seq_id = ubatch.seq_id_unq[s];
const uint32_t rollback = seq_id >= 0 && (uint32_t) seq_id < rs_idx.size() ? rs_idx[seq_id] : 0;
LLAMA_LOG_DEBUG("%s: seq %d pos [%d, %d] rollback=%u\n", __func__, seq_id,
ubatch.pos[0], ubatch.pos[ubatch.n_tokens - 1], rollback);
}
}
return plan;
@@ -704,8 +748,17 @@ static std::vector<llama_kv_cache_dsv4_context::comp_plan> dsv4_build_comp_plans
std::vector<llama_kv_cache_dsv4_context::comp_plan> plans;
plans.reserve(ubatches.size());
// the first ubatch touching a seq consumes its rollback restore
std::vector<uint32_t> rs(rs_idx);
for (const llama_ubatch & ubatch : ubatches) {
plans.push_back(dsv4_build_comp_plan(ubatch, ratio, overlap, state_size, kv_size, n_stream, n_rs_seq, rs_idx));
plans.push_back(dsv4_build_comp_plan(ubatch, ratio, overlap, state_size, kv_size, n_stream, n_rs_seq, rs));
for (uint32_t s = 0; s < ubatch.n_seqs_unq; ++s) {
const llama_seq_id seq_id = ubatch.seq_id_unq[s];
if (seq_id >= 0 && (size_t) seq_id < rs.size()) {
rs[seq_id] = 0;
}
}
}
return plans;
@@ -803,16 +856,15 @@ static llama_kv_cache_dsv4_context::comp_plan dsv4_build_reserve_comp_plan(
return plan;
}
const uint32_t n_seqs = std::max<uint32_t>(1, ubatch.n_seqs);
const uint32_t n_seq_tokens = std::max<uint32_t>(1, ubatch.n_seq_tokens);
const uint64_t n_blocks_u64 = (uint64_t) n_seqs*((n_seq_tokens + ratio - 1)/ratio);
const size_t n_blocks = (size_t) std::max<uint64_t>(1, n_blocks_u64);
GGML_ASSERT((uint64_t) n_blocks == std::max<uint64_t>(1, n_blocks_u64));
// worst case over every seq split: sum of per-seq ceil(tokens/ratio) is at
// most floor(n_tokens/ratio) + n_seqs
const uint32_t n_seqs = std::max<uint32_t>(1, ubatch.n_seqs);
const size_t n_blocks = (size_t) ubatch.n_tokens/ratio + n_seqs;
const uint64_t state_rows = (uint64_t) state_size*n_stream;
const size_t n_persist = (size_t) std::min<uint64_t>(ubatch.n_tokens, state_rows);
const size_t n_restore = n_rs_seq > 0 ? (size_t) state_size*std::max<uint32_t>(1, ubatch.n_seqs_unq) : 0;
const size_t n_snapshot = (size_t) n_rs_seq*state_size*std::max<uint32_t>(1, ubatch.n_seqs_unq);
const size_t n_restore = n_rs_seq > 0 ? (size_t) state_size*n_stream : 0;
const size_t n_snapshot = (size_t) n_rs_seq*state_size*n_stream;
plan.state_pos .resize(ubatch.n_tokens);
plan.state_persist_src_idxs.resize(n_persist);
@@ -1356,7 +1408,9 @@ llama_memory_context_ptr llama_kv_cache_dsv4::init_batch(
if (has_coupled) {
ubatch = balloc.split_seq(n_ubatch);
} else {
ubatch = balloc.split_equal(n_ubatch, raw_per_seq || comp_per_seq, 0);
// [TAG_RECURRENT_ROLLBACK_SPLITS]
// the trailing (1 + n_rs_seq) tokens of each seq must stay in the same ubatch
ubatch = balloc.split_equal(n_ubatch, raw_per_seq || comp_per_seq, n_rs_seq > 0 ? n_rs_seq + 1 : 0);
}
if (ubatch.n_tokens == 0) {
@@ -1433,6 +1487,11 @@ bool llama_kv_cache_dsv4::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_pos p1
return false;
}
// pending rollback is single-use: stacked partial removals don't compose
if (rs_idx[seq_id] != 0) {
return false;
}
const bool res = kv_raw->seq_rm(seq_id, p0, p1);
if (res) {
rs_idx[seq_id] = (uint32_t) rollback;
@@ -1594,9 +1653,7 @@ void llama_kv_cache_dsv4::state_read(llama_io_read_i & io, llama_seq_id seq_id,
kv_raw->state_read(io, seq_id, flags);
if (!partial_only) {
kv_csa->clear(true);
kv_hca->clear(true);
kv_lid->clear(true);
clear_compressed(seq_id, true);
dsv4_state_read_k_cache(io, kv_csa.get(), seq_id, flags);
dsv4_state_read_k_cache(io, kv_hca.get(), seq_id, flags);
@@ -1680,6 +1737,7 @@ void llama_kv_cache_dsv4::clear_compressed(llama_seq_id seq_id, bool data) {
kv->seq_rm(seq_id, -1, -1);
if (data) {
//TODO: do not clear the kv-cache during `seq_rm`, ref: https://github.com/ggml-org/llama.cpp/pull/26490#discussion_r3798143663
for (uint32_t il : kv->get_layer_ids()) {
dsv4_clear_tensor_stream(kv->get_k_storage(il), (uint32_t) seq_id);
}
src/llama-kv-cache.cpp
+2
View File
@@ -383,6 +383,7 @@ bool llama_kv_cache::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_pos p1) {
return true;
}
// TODO: fix incosistent handling of `seq_id < 0` and `seq_id == -1` in the codebase [TAG_LLAMA_SEQ_ID_NEG]
GGML_ASSERT(seq_id == -1 || (seq_id >= 0 && (size_t) seq_id < seq_to_stream.size()));
if (p0 < 0) {
@@ -2048,6 +2049,7 @@ void llama_kv_cache::state_read(llama_io_read_i & io, llama_seq_id seq_id, llama
GGML_UNUSED(flags);
// TODO: fix incosistent handling of `seq_id < 0` and `seq_id == -1` in the codebase [TAG_LLAMA_SEQ_ID_NEG]
GGML_ASSERT(seq_id == -1 || (seq_id >= 0 && (size_t) seq_id < seq_to_stream.size()));
uint32_t n_stream_cur;
src/llama-memory-recurrent.cpp
+21 -13
View File
@@ -158,13 +158,14 @@ bool llama_memory_recurrent::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_pos
p1 = std::numeric_limits<llama_pos>::max();
}
if ((uint32_t) seq_id >= this->n_seq_max) {
LLAMA_LOG_ERROR("%s: invalid seq_id (%d) - larger than n_seq_max (%d)\n", __func__, seq_id, this->n_seq_max);
return false;
}
const bool rm_all = p0 == 0 && p1 == std::numeric_limits<llama_pos>::max();
if (rm_all) {
if (seq_id >= 0) {
set_rs_idx(seq_id, 0);
} else {
std::fill(rs_idx.begin(), rs_idx.end(), 0);
}
set_rs_idx(seq_id, 0);
}
// models like Mamba or RWKV can't have a state partially erased at the end
@@ -181,7 +182,9 @@ bool llama_memory_recurrent::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_pos
// partial rollback via per-token snapshot index (bounded by n_rs_seq)
if (0 < p0 && p0 <= cell.pos && p1 > cell.pos) {
const llama_pos rollback = cell.pos - (p0 - 1);
if (rollback >= 1 && rollback <= (llama_pos) n_rs_seq) {
// pending rollback is single-use
const bool pending = rs_idx[seq_id] != 0;
if (!pending && rollback >= 1 && rollback <= (llama_pos) n_rs_seq) {
set_rs_idx(seq_id, (uint32_t) rollback);
cell.pos = p0 - 1;
return true;
@@ -390,10 +393,17 @@ llama_pos llama_memory_recurrent::seq_pos_max(llama_seq_id seq_id) const {
}
void llama_memory_recurrent::set_rs_idx(llama_seq_id seq_id, uint32_t idx) {
if (seq_id < 0 || (size_t) seq_id >= rs_idx.size()) {
if (seq_id < 0) {
std::fill(rs_idx.begin(), rs_idx.end(), 0);
return;
}
rs_idx[seq_id] = (idx > n_rs_seq) ? n_rs_seq : idx;
assert(n_seq_max == rs_idx.size());
GGML_ASSERT((uint32_t) seq_id < n_seq_max);
GGML_ASSERT(idx <= n_rs_seq);
rs_idx[seq_id] = idx;
}
std::map<ggml_backend_buffer_type_t, size_t> llama_memory_recurrent::memory_breakdown() const {
@@ -742,6 +752,7 @@ void llama_memory_recurrent::state_write(llama_io_write_i & io, llama_seq_id seq
uint32_t cell_range_begin = size;
for (uint32_t i = 0; i < size; ++i) {
const auto & cell = cells[i];
// TODO: fix incosistent handling of `seq_id < 0` and `seq_id == -1` in the codebase [TAG_LLAMA_SEQ_ID_NEG]
if ((seq_id == -1 && !cell.is_empty()) || cell.has_seq_id(seq_id)) {
++cell_count;
uint32_t rs_idx_cur = 0;
@@ -827,6 +838,7 @@ void llama_memory_recurrent::state_read(llama_io_read_i & io, llama_seq_id seq_i
}
if (!res) {
// TODO: fix incosistent handling of `seq_id < 0` and `seq_id == -1` in the codebase [TAG_LLAMA_SEQ_ID_NEG]
if (seq_id == -1) {
clear(true);
} else {
@@ -836,11 +848,7 @@ void llama_memory_recurrent::state_read(llama_io_read_i & io, llama_seq_id seq_i
}
if (n_rs_seq != 0) {
if (seq_id == -1) {
std::fill(rs_idx.begin(), rs_idx.end(), 0);
} else {
set_rs_idx(seq_id, 0);
}
set_rs_idx(seq_id, 0);
}
}
src/llama-model-saver.cpp
+20
View File
@@ -293,6 +293,21 @@ void llama_model_saver::add_kv_from_model() {
add_kv(LLM_KV_ATTENTION_INDEXER_LOCAL_BLOCKS, hparams.indexer_local_blocks);
add_kv(LLM_KV_ATTENTION_INDEXER_TYPES, hparams.is_indexer_full_impl, true);
add_kv(LLM_KV_ATTENTION_RECURRENT_LAYERS, hparams.is_recr_impl, true);
add_kv(LLM_KV_ATTENTION_OUTPUT_GROUP_COUNT, hparams.dsv4_o_group_count);
add_kv(LLM_KV_ATTENTION_OUTPUT_LORA_RANK, hparams.dsv4_o_lora_rank);
add_kv(LLM_KV_ATTENTION_COMPRESS_ROPE_FREQ_BASE, hparams.dsv4_compress_rope_base);
if (model->arch == LLM_ARCH_DEEPSEEK4 || hparams.dsv4_hc_mult > 0) {
// the loader requires one compress ratio per layer, including nextn layers
const std::vector<uint32_t> compress_ratios(
hparams.dsv4_compress_ratios.begin(), hparams.dsv4_compress_ratios.begin() + hparams.n_layer_all);
add_kv(LLM_KV_ATTENTION_COMPRESS_RATIOS, compress_ratios);
} else {
add_kv(LLM_KV_ATTENTION_COMPRESS_RATIOS, hparams.dsv4_compress_ratios, true);
}
add_kv(LLM_KV_HYPER_CONNECTION_COUNT, hparams.dsv4_hc_mult);
add_kv(LLM_KV_HYPER_CONNECTION_SINKHORN_ITERATIONS, hparams.dsv4_hc_sinkhorn_iters);
add_kv(LLM_KV_HYPER_CONNECTION_EPSILON, hparams.dsv4_hc_eps);
add_kv(LLM_KV_HASH_LAYER_COUNT, hparams.dsv4_hash_layer_count);
const float rope_scaling_factor = hparams.rope_freq_scale_train == 1.0f ? 0.0f : 1.0f/hparams.rope_freq_scale_train;
@@ -417,11 +432,16 @@ void llama_model_saver::add_tensors_from_model() {
add_tensor(model->output_s);
add_tensor(model->output_in_s);
add_tensor(model->output_res_score);
add_tensor(model->nextn_proj_pre);
add_tensor(model->nextn_proj_post);
add_tensor(model->cls);
add_tensor(model->cls_b);
add_tensor(model->cls_out);
add_tensor(model->cls_out_b);
add_tensor(model->cls_norm);
add_tensor(model->hc_head_fn);
add_tensor(model->hc_head_base);
add_tensor(model->hc_head_scale);
for (const struct llama_layer & layer : model->layers) {
for (size_t i = 0; i < sizeof(layer)/sizeof(struct ggml_tensor *); ++i) {
src/llama-model.cpp
+90 -6
View File
@@ -516,6 +516,8 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
const llama_meta_device_get_split_state_userdata * ud = (const llama_meta_device_get_split_state_userdata *) userdata;
const llama_hparams & hparams = ud->model->hparams;
const std::string tensor_name = tensor->name;
const bool is_dsv4 = ud->model->arch == LLM_ARCH_DEEPSEEK4 ||
(ud->model->arch == LLM_ARCH_DFLASH && hparams.dsv4_hc_mult > 0);
static const std::regex pattern_q_weight ("blk\\.\\d*\\.attn_q.weight");
static const std::regex pattern_kv_weight ("blk\\.\\d*\\.attn_(k|v).weight");
@@ -525,9 +527,13 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
static const std::regex pattern_qkv_bias ("blk\\.\\d*\\.attn_qkv.bias");
static const std::regex pattern_qk_norm ("blk\\.\\d*\\.attn_(q|k)_norm\\.weight");
static const std::regex pattern_kv_cache ("cache_(k|v)_l\\d*");
static const std::regex pattern_dsv4_state ("dsv4_(csa|hca|lid)_state_(kv|score)_l\\d*");
static const std::regex pattern_attn_sinks ("blk\\.\\d*\\.attn_sinks.weight");
static const std::regex pattern_attn_out_weight ("blk\\.\\d*\\.attn_output.weight");
static const std::regex pattern_attn_out_bias ("blk\\.\\d*\\.attn_output.bias");
static const std::regex pattern_attn_out_a_weight("blk\\.\\d*\\.attn_output_a\\.weight");
static const std::regex pattern_attn_out_b_weight("blk\\.\\d*\\.attn_output_b\\.weight");
static const std::regex pattern_attn_q_b_weight ("blk\\.\\d*\\.attn_q_b\\.weight");
static const std::regex pattern_attn_gate_weight("blk\\.\\d*\\.attn_gate.weight");
static const std::regex pattern_ssm_dt ("blk\\.\\d*\\.ssm_dt.bias");
@@ -546,8 +552,11 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
static const std::regex pattern_ffn_gate_bias ("blk\\.\\d*\\.ffn_gate(_exps)?.bias");
static const std::regex pattern_ffn_gate_up_weight("blk\\.\\d*\\.ffn_gate_up(_exps)?.weight");
static const std::regex pattern_ffn_down_weight ("blk\\.\\d*\\.ffn_down(_exps)?.weight");
static const std::regex pattern_ffn_down_bias ("blk\\.\\d*\\.ffn_down.bias");
static const std::regex pattern_ffn_down_exps_bias("blk\\.\\d*\\.ffn_down_exps.bias");
static const std::regex pattern_ffn_down_bias ("blk\\.\\d*\\.ffn_down.bias");
static const std::regex pattern_ffn_down_exps_bias ("blk\\.\\d*\\.ffn_down_exps.bias");
static const std::regex pattern_ffn_up_shexp_weight ("blk\\.\\d*\\.ffn_up_shexp.weight");
static const std::regex pattern_ffn_gate_shexp_weight ("blk\\.\\d*\\.ffn_gate_shexp.weight");
static const std::regex pattern_ffn_down_shexp_weight ("blk\\.\\d*\\.ffn_down_shexp.weight");
static const std::regex pattern_output_weight("output\\.weight");
static const std::regex pattern_output_bias ("output\\.bias");
@@ -604,6 +613,32 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
};
auto get_tensor_config = [&]() -> tensor_config {
if (is_dsv4) {
if (std::regex_match(tensor_name, pattern_kv_cache) ||
std::regex_match(tensor_name, pattern_dsv4_state)) {
return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_MIRRORED);
}
if (std::regex_match(tensor_name, pattern_attn_sinks)) {
return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_0, "attn_output_a.weight");
}
if (std::regex_match(tensor_name, pattern_attn_q_b_weight)) {
return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_1, "attn_output_a.weight");
}
if (std::regex_match(tensor_name, pattern_attn_out_a_weight)) {
return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_2);
}
if (std::regex_match(tensor_name, pattern_attn_out_b_weight)) {
return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_0);
}
if (std::regex_match(tensor_name, pattern_ffn_up_shexp_weight) ||
std::regex_match(tensor_name, pattern_ffn_gate_shexp_weight)) {
return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_1, "ffn_down_shexp.weight");
}
if (std::regex_match(tensor_name, pattern_ffn_down_shexp_weight)) {
return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_0, "ffn_down_shexp.weight");
}
}
// standard attention
if (std::regex_match(tensor_name, pattern_q_weight) || std::regex_match(tensor_name, pattern_kv_weight)) {
return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_1, "attn_output.weight", "ssm_out.weight");
@@ -671,11 +706,14 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_MIRRORED);
}
if (std::regex_match(tensor_name, pattern_ffn_down_exps_bias)) {
return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_PARTIAL);
return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_PARTIAL, "ffn_down_exps.weight");
}
// output
if (std::regex_match(tensor_name, pattern_output_weight)) {
if (is_dsv4) {
return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_MIRRORED);
}
return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_1);
}
if (std::regex_match(tensor_name, pattern_output_bias)) {
@@ -705,6 +743,9 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
GGML_ASSERT(tensor->ne[axis] == 2*key_dim + value_dim);
return {{key_dim, 2}, {value_dim, 1}};
}
if (std::regex_match(tensor_name, pattern_r_cache)) {
return {{key_dim * (hparams.ssm_d_conv - 1), 2}, {value_dim * (hparams.ssm_d_conv - 1), 1}};
}
} else {
const int64_t head_ratio = n_v_heads / n_k_heads;
if (std::regex_match(tensor_name, pattern_qkv_weight) || std::regex_match(tensor_name, pattern_ssm_conv1d)) {
@@ -793,12 +834,34 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
blck_size_perf *= 2;
}
const int64_t granularity_q = std::lcm(n_embd_q, blck_size_perf);
const int64_t granularity_head = granularity_q / hparams.n_embd_head_k(il); // for tensors with one value per head
if (std::regex_match(tensor_name, pattern_attn_sinks)) {
GGML_ASSERT(segments.size() == 1);
return {std::lcm(n_embd_q, blck_size_perf)/n_embd_q * n_gqa};
if (is_dsv4) {
return {hparams.n_head(il) / hparams.dsv4_o_group_count};
}
return {granularity_head};
}
const int64_t granularity_q = std::lcm(n_embd_q, blck_size_perf);
if (is_dsv4) {
if (std::regex_match(tensor_name, pattern_attn_q_b_weight)) {
GGML_ASSERT(segments.size() == 1);
// the grouped output projection requires each device to hold whole groups of heads
const int64_t n_head_group = hparams.n_head(il) / hparams.dsv4_o_group_count;
return {n_head_group * hparams.n_embd_head_k(il)};
}
if (std::regex_match(tensor_name, pattern_attn_out_a_weight)) {
GGML_ASSERT(segments.size() == 1);
return {1};
}
if (std::regex_match(tensor_name, pattern_attn_out_b_weight)) {
GGML_ASSERT(segments.size() == 1);
// the boundaries must align with wo_a's per-group split, so quant blocks must not straddle groups
GGML_ASSERT(hparams.dsv4_o_lora_rank % blck_size == 0);
return {hparams.dsv4_o_lora_rank};
}
}
if (std::regex_match(tensor_name, pattern_q_weight) || std::regex_match(tensor_name, pattern_q_bias)) {
GGML_ASSERT(segments.size() == 1);
// some models have Q gate tensors, for those cases the granularity needs to be doubled:
@@ -811,6 +874,13 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
GGML_ASSERT(segments.size() == 1);
return {granularity_q};
}
if (std::regex_match(tensor_name, pattern_attn_gate_weight)) {
GGML_ASSERT(segments.size() == 1);
if (tensor->ne[1] == hparams.n_head(il)) {
return {granularity_head};
}
return {granularity_q};
}
const int64_t granularity_kv = granularity_q / n_gqa;
if (std::regex_match(tensor_name, pattern_kv_weight) ||
@@ -828,7 +898,11 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
// FFN
if (std::regex_match(tensor_name, pattern_ffn_up_weight) || std::regex_match(tensor_name, pattern_ffn_up_bias) ||
std::regex_match(tensor_name, pattern_ffn_gate_weight) || std::regex_match(tensor_name, pattern_ffn_gate_bias) ||
std::regex_match(tensor_name, pattern_ffn_gate_up_weight) || std::regex_match(tensor_name, pattern_ffn_down_weight)) {
std::regex_match(tensor_name, pattern_ffn_gate_up_weight) ||
std::regex_match(tensor_name, pattern_ffn_down_weight) ||
std::regex_match(tensor_name, pattern_ffn_up_shexp_weight) ||
std::regex_match(tensor_name, pattern_ffn_gate_shexp_weight) ||
std::regex_match(tensor_name, pattern_ffn_down_shexp_weight)) {
const int64_t blck_size_perf = std::lcm(blck_size, 128);
GGML_ASSERT(segments.size() == 1);
return {blck_size_perf};
@@ -879,6 +953,16 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
memset(split_state.ne, 0, sizeof(split_state.ne));
split_state.nr[0] = 1;
split_state.n_segments = 1;
if (split_state.axis == GGML_BACKEND_SPLIT_AXIS_PARTIAL) {
GGML_ASSERT(tc.tensor_axis_0 != tensor);
const ggml_backend_meta_split_state source_split_state = llama_meta_device_get_split_state(tc.tensor_axis_0, userdata);
GGML_ASSERT(source_split_state.axis >= 0 && source_split_state.axis < GGML_MAX_DIMS);
for (size_t j = 0; j < ud->n_devices; j++) {
for (size_t is = 0; is < source_split_state.n_segments; is++) {
split_state.ne[j] += source_split_state.ne[is*ud->n_devices + j] * source_split_state.nr[is];
}
}
}
}
return split_state;
GGML_UNUSED(userdata);
src/models/deepseek4.cpp
+2
View File
@@ -1,3 +1,4 @@
#include "llama-hparams.h"
#include "models.h"
#include "llama-kv-cache-dsv4.h"
@@ -58,6 +59,7 @@ void llama_model_deepseek4::load_arch_hparams(llama_model_loader & ml) {
if (n_compress_ratios < hparams.n_layer_all) {
throw std::runtime_error("DeepSeek-V4 compress_ratios is shorter than block_count");
}
GGML_ASSERT(n_compress_ratios <= LLAMA_MAX_LAYERS);
ml.get_arr(LLM_KV_ATTENTION_COMPRESS_RATIOS, hparams.dsv4_compress_ratios);
ml.get_key(LLM_KV_EXPERT_GATING_FUNC, hparams.expert_gating_func);
src/models/dflash.cpp
+4 -3
View File
@@ -117,6 +117,10 @@ void llama_model_dflash::load_arch_tensors(llama_model_loader &) {
output_norm_enc = create_tensor(tn(LLM_TENSOR_ENC_OUTPUT_NORM, "weight"), { n_embd }, 0); // encoder hidden_norm (after fc)
output_norm = create_tensor(tn(LLM_TENSOR_OUTPUT_NORM, "weight"), { n_embd }, 0); // decoder final norm
// optional: reduced-vocab drafts ship their own lm head, full-vocab drafts can share the target's via ctx_other
// a draft with its own embeddings + head references no target tensors and can run on devices the target does not use (e.g. -devd with a tensor-split target)
output = create_tensor(tn(LLM_TENSOR_OUTPUT, "weight"), { n_embd, n_vocab_draft }, TENSOR_NOT_REQUIRED);
if (hparams.dsv4_hc_mult > 0) {
const int64_t q_lora_rank = hparams.n_lora_q;
const int64_t n_ff_exp = hparams.n_ff_exp;
@@ -167,9 +171,6 @@ void llama_model_dflash::load_arch_tensors(llama_model_loader &) {
return;
}
// optional: reduced-vocab drafts ship their own, full-vocab drafts share the target's via ctx_other
output = create_tensor(tn(LLM_TENSOR_OUTPUT, "weight"), { n_embd, n_vocab_draft }, TENSOR_NOT_REQUIRED);
for (int i = 0; i < n_layer; ++i) {
auto & layer = layers[i];
src/models/glm4-moe.cpp
+186 -16
View File
@@ -29,10 +29,19 @@ void llama_model_glm4_moe::load_arch_hparams(llama_model_loader & ml) {
}
}
void llama_model_glm4_moe::load_arch_tensors(llama_model_loader &) {
void llama_model_glm4_moe::load_arch_tensors(llama_model_loader & ml) {
LLAMA_LOAD_LOCALS;
const int64_t n_expert_shared = hparams.n_expert_shared;
const bool mtp_only = (hparams.n_layer_nextn > 0) && (ml.get_weight("blk.0.attn_norm.weight") == nullptr);
const std::string mtp_probe = "blk." + std::to_string(n_layer) + ".nextn.eh_proj.weight";
const bool trunk_only = (hparams.n_layer_nextn > 0) && (ml.get_weight(mtp_probe.c_str()) == nullptr);
const int trunk_flags = mtp_only ? TENSOR_NOT_REQUIRED : 0;
int mtp_flags = trunk_only ? TENSOR_NOT_REQUIRED : 0;
if (!ml.load_mtp) {
mtp_flags |= TENSOR_SKIP;
}
GGML_ASSERT(hparams.n_expert > 0 && "n_expert must be > 0 for GLM4_MOE MoE layers");
GGML_ASSERT(hparams.n_expert_used > 0 && "n_expert_used must be > 0 for GLM4_MOE MoE layers");
@@ -47,16 +56,9 @@ void llama_model_glm4_moe::load_arch_tensors(llama_model_loader &) {
output = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), { n_embd, n_vocab }, TENSOR_DUPLICATED);
}
// Load ALL tensors including NextN layer to satisfy total tensor count
// but only PROCESS up to last layer (skipping final NextN layer) in forward pass
for (int i = 0; i < n_layer_all; ++i) {
int flags = 0;
if (i >= n_layer) {
// skip all tensors in the NextN layers
flags |= TENSOR_SKIP;
}
auto & layer = layers[i];
const int flags = i < n_layer ? trunk_flags : mtp_flags;
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), { n_embd }, flags);
@@ -110,24 +112,186 @@ void llama_model_glm4_moe::load_arch_tensors(llama_model_loader &) {
layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), { n_embd, n_ff }, flags);
}
// NextN/MTP tensors (preserved but unused) - conditionally load for last nextn_predict_layers
// NextN/MTP tensors
if (i >= n_layer) {
layer.nextn.eh_proj = create_tensor(tn(LLM_TENSOR_NEXTN_EH_PROJ, "weight", i), { 2 * n_embd, n_embd }, flags);
layer.nextn.enorm = create_tensor(tn(LLM_TENSOR_NEXTN_ENORM, "weight", i), { n_embd }, flags);
layer.nextn.hnorm = create_tensor(tn(LLM_TENSOR_NEXTN_HNORM, "weight", i), { n_embd }, flags);
// Optional tensors
layer.nextn.embed_tokens = create_tensor(tn(LLM_TENSOR_NEXTN_EMBED_TOKENS, "weight", i), { n_embd, n_vocab }, flags | TENSOR_NOT_REQUIRED);
layer.nextn.shared_head_head = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_HEAD, "weight", i), { n_embd, n_vocab }, flags | TENSOR_NOT_REQUIRED);
layer.nextn.shared_head_norm = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_NORM, "weight", i), { n_embd }, flags | TENSOR_NOT_REQUIRED);
layer.nextn.embed_tokens = create_tensor(tn(LLM_TENSOR_NEXTN_EMBED_TOKENS, "weight", i), { n_embd, n_vocab }, TENSOR_NOT_REQUIRED | flags);
layer.nextn.shared_head_head = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_HEAD, "weight", i), { n_embd, n_vocab }, TENSOR_NOT_REQUIRED | flags);
layer.nextn.shared_head_norm = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_NORM, "weight", i), { n_embd }, TENSOR_NOT_REQUIRED | flags);
}
}
}
std::unique_ptr<llm_graph_context> llama_model_glm4_moe::build_arch_graph(const llm_graph_params & params) const {
if (params.gtype == LLM_GRAPH_TYPE_DECODER_MTP) {
return std::make_unique<graph_mtp>(*this, params);
}
return std::make_unique<graph>(*this, params);
}
llama_model_glm4_moe::graph_mtp::graph_mtp(const llama_model & model, const llm_graph_params & params)
: llm_graph_context(params) {
GGML_ASSERT(hparams.n_layer_nextn > 0 && "GLM4_MOE MTP requires n_layer_nextn > 0");
GGML_ASSERT(hparams.n_layer_nextn == 1 && "GLM4_MOE MTP currently only supports a single MTP block");
const int64_t n_embd_head = hparams.n_embd_head_v();
GGML_ASSERT(n_embd_head == hparams.n_embd_head_k());
const int il = hparams.n_layer() + cparams.nextn_layer_offset;
GGML_ASSERT(cparams.nextn_layer_offset >= 0 &&
cparams.nextn_layer_offset < (int) hparams.n_layer_nextn &&
"nextn_layer_offset out of range [0, n_layer_nextn)");
const auto & layer = model.layers[il];
GGML_ASSERT(layer.nextn.eh_proj && "MTP block missing nextn.eh_proj");
GGML_ASSERT(layer.nextn.enorm && "MTP block missing nextn.enorm");
GGML_ASSERT(layer.nextn.hnorm && "MTP block missing nextn.hnorm");
GGML_ASSERT(layer.ffn_gate_inp && "MTP block missing ffn_gate_inp");
auto inp = std::make_unique<llm_graph_input_embd_h>(hparams.n_embd);
inp->tokens = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, n_tokens);
ggml_set_input(inp->tokens);
inp->embd = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, hparams.n_embd_inp(), n_tokens);
ggml_set_input(inp->embd);
ggml_tensor * tok_embd;
if (ubatch.token) {
ggml_tensor * tok_embd_w = layer.nextn.embed_tokens ? layer.nextn.embed_tokens : model.tok_embd;
tok_embd = ggml_get_rows(ctx0, tok_embd_w, inp->tokens);
} else {
tok_embd = inp->embd;
}
cb(tok_embd, "mtp_tok_embd", il);
inp->h = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, hparams.n_embd, n_tokens);
ggml_set_input(inp->h);
ggml_set_name(inp->h, "mtp_h_input");
ggml_tensor * h_embd = inp->h;
res->add_input(std::move(inp));
ggml_tensor * inp_pos = build_inp_pos();
ggml_tensor * inp_out_ids = build_inp_out_ids();
auto * inp_attn = build_attn_inp_kv();
ggml_tensor * h_norm = build_norm(h_embd, layer.nextn.hnorm, nullptr, LLM_NORM_RMS, il);
cb(h_norm, "mtp_hnorm", il);
ggml_tensor * e_norm = build_norm(tok_embd, layer.nextn.enorm, nullptr, LLM_NORM_RMS, il);
cb(e_norm, "mtp_enorm", il);
ggml_tensor * concat = ggml_concat(ctx0, e_norm, h_norm, 0);
cb(concat, "mtp_concat", il);
ggml_tensor * cur = build_lora_mm(layer.nextn.eh_proj, concat, layer.nextn.eh_proj_s);
cb(cur, "mtp_eh_proj", il);
ggml_tensor * inpSA = cur;
cur = build_norm(cur, layer.attn_norm, nullptr, LLM_NORM_RMS, il);
cb(cur, "mtp_attn_norm", il);
auto [Qcur, Kcur, Vcur] = build_qkv(layer, cur,
n_embd_head, n_head, n_head_kv, il);
if (layer.attn_q_norm) {
Qcur = build_norm(Qcur, layer.attn_q_norm, nullptr, LLM_NORM_RMS, il);
cb(Qcur, "mtp_Qcur_normed", il);
}
if (layer.attn_k_norm) {
Kcur = build_norm(Kcur, layer.attn_k_norm, nullptr, LLM_NORM_RMS, il);
cb(Kcur, "mtp_Kcur_normed", il);
}
Qcur = ggml_rope_ext(ctx0, Qcur, inp_pos, nullptr, n_rot,
rope_type, n_ctx_orig, freq_base, freq_scale,
ext_factor, attn_factor, beta_fast, beta_slow);
Kcur = ggml_rope_ext(ctx0, Kcur, inp_pos, nullptr, n_rot,
rope_type, n_ctx_orig, freq_base, freq_scale,
ext_factor, attn_factor, beta_fast, beta_slow);
cb(Qcur, "mtp_Qcur", il);
cb(Kcur, "mtp_Kcur", il);
cb(Vcur, "mtp_Vcur", il);
cur = build_attn(inp_attn,
layer.wo, nullptr, layer.wo_s,
Qcur, Kcur, Vcur, nullptr, nullptr, nullptr,
1.0f / sqrtf(float(n_embd_head)), il);
cb(cur, "mtp_attn_out", il);
ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpSA);
cb(ffn_inp, "mtp_ffn_inp", il);
cur = build_norm(ffn_inp, layer.attn_post_norm, nullptr, LLM_NORM_RMS, il);
cb(cur, "mtp_post_attn_norm", il);
ggml_tensor * routed_out = build_moe_ffn(cur,
layer.ffn_gate_inp,
layer.ffn_up_exps,
layer.ffn_gate_exps,
layer.ffn_down_exps,
layer.ffn_exp_probs_b,
n_expert, n_expert_used,
LLM_FFN_SILU, hparams.expert_weights_norm,
hparams.expert_weights_scale,
(llama_expert_gating_func_type) hparams.expert_gating_func,
il);
cb(routed_out, "mtp_ffn_moe_out", il);
ggml_tensor * shared_out = build_ffn(cur,
layer.ffn_up_shexp, nullptr, nullptr,
layer.ffn_gate_shexp, nullptr, nullptr,
layer.ffn_down_shexp, nullptr, nullptr,
nullptr,
LLM_FFN_SILU, LLM_FFN_PAR, il);
cb(shared_out, "mtp_ffn_shexp_out", il);
cur = ggml_add(ctx0, routed_out, shared_out);
cb(cur, "mtp_ffn_out", il);
cur = ggml_add(ctx0, cur, ffn_inp);
cb(cur, "mtp_post_ffn", il);
ggml_tensor * head_norm_w = layer.nextn.shared_head_norm
? layer.nextn.shared_head_norm
: model.output_norm;
GGML_ASSERT(head_norm_w && "GLM4_MOE MTP: missing both nextn.shared_head_norm and output_norm");
cur = build_norm(cur, head_norm_w, nullptr, LLM_NORM_RMS, -1);
cb(cur, "h_nextn", -1);
res->t_h_nextn = cur;
if (inp_out_ids) {
cur = ggml_get_rows(ctx0, cur, inp_out_ids);
}
cb(cur, "mtp_shared_head_norm", -1);
ggml_tensor * head_w = layer.nextn.shared_head_head
? layer.nextn.shared_head_head
: model.output;
ggml_tensor * head_s = layer.nextn.shared_head_head
? layer.nextn.shared_head_head_s
: model.output_s;
GGML_ASSERT(head_w && "GLM4_MOE MTP: missing LM head (nextn.shared_head_head or model.output)");
cur = build_lora_mm(head_w, cur, head_s);
cb(cur, "result_output", -1);
res->t_logits = cur;
ggml_build_forward_expand(gf, cur);
}
llama_model_glm4_moe::graph::graph(const llama_model & model, const llm_graph_params & params) : llm_graph_context(params) {
const int64_t n_embd_head = hparams.n_embd_head_v();
@@ -154,8 +318,7 @@ llama_model_glm4_moe::graph::graph(const llama_model & model, const llm_graph_pa
ggml_tensor * inp_out_ids = build_inp_out_ids();
// Only process up to last layer (skip final NextN layer)
// Final layer tensors are loaded but not processed in forward pass
// NextN layers are processed by graph_mtp.
for (int il = 0; il < n_layer; ++il) {
ggml_tensor * inpSA = inpL;
@@ -205,7 +368,7 @@ llama_model_glm4_moe::graph::graph(const llama_model & model, const llm_graph_pa
model.layers[il].wo, NULL, model.layers[il].wo_s,
Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, 1.0f/sqrtf(float(n_embd_head)), il);
}
if (il == n_layer - 1 && inp_out_ids) {
if (il == n_layer - 1 && inp_out_ids && (!cparams.embeddings_nextn || cparams.embeddings_nextn_masked)) {
cur = ggml_get_rows(ctx0, cur, inp_out_ids);
inpSA = ggml_get_rows(ctx0, inpSA, inp_out_ids);
}
@@ -265,6 +428,13 @@ llama_model_glm4_moe::graph::graph(const llama_model & model, const llm_graph_pa
cur = inpL;
cur = build_norm(cur, model.output_norm, NULL, LLM_NORM_RMS, -1);
cb(cur, "h_nextn", -1);
res->t_h_nextn = cur;
if (cparams.embeddings_nextn && !cparams.embeddings_nextn_masked && inp_out_ids) {
cur = ggml_get_rows(ctx0, cur, inp_out_ids);
}
cb(cur, "result_norm", -1);
res->t_embd = cur;
src/models/mamba-base.cpp
+7 -9
View File
@@ -182,13 +182,14 @@ ggml_tensor * llm_build_mamba_base::build_mamba2_layer(llm_graph_input_rs * inp,
ggml_tensor * conv = build_rs(inp, conv_states_all, hparams.n_embd_r(), n_seqs);
conv = ggml_reshape_3d(ctx0, conv, d_conv - 1, d_inner + 2 * n_group * d_state, n_seqs);
// {n_embd, n_tokens} => {n_embd, n_seq_tokens, n_seqs}
cur = ggml_reshape_3d(ctx0, cur, cur->ne[0], n_seq_tokens, n_seqs);
// d_in_proj = 2 * self.d_inner + 2 * self.ngroups * self.d_state + self.nheads
// {n_embd, d_in_proj} @ {n_embd, n_seq_tokens, n_seqs} => {d_in_proj, n_seq_tokens, n_seqs}
// Keep the projection 2D: with a {n_embd, 1, n_seqs} batch the CUDA backend
// dispatches a column-batched GEMV for what is a large dense GEMM.
// {n_embd, d_in_proj} @ {n_embd, n_tokens} => {d_in_proj, n_tokens}
ggml_tensor * zxBCdt = build_lora_mm(model.layers[il].ssm_in, cur, model.layers[il].ssm_in_s);
// {d_in_proj, n_tokens} => {d_in_proj, n_seq_tokens, n_seqs}
zxBCdt = ggml_reshape_3d(ctx0, zxBCdt, zxBCdt->ne[0], n_seq_tokens, n_seqs);
// split the above in three
ggml_tensor * z = ggml_view_4d(ctx0, zxBCdt, head_dim, n_head, n_seq_tokens, n_seqs, head_dim * zxBCdt->nb[0],
@@ -290,15 +291,12 @@ ggml_tensor * llm_build_mamba_base::build_mamba2_layer(llm_graph_input_rs * inp,
y = build_norm(y, model.layers[il].ssm_norm, NULL, LLM_NORM_RMS, il);
}
y = ggml_reshape_3d(ctx0, y, d_inner, n_seq_tokens, n_seqs);
y = ggml_reshape_2d(ctx0, y, d_inner, n_seq_tokens * n_seqs);
// {d_inner, n_embd} @ {d_inner, n_seq_tokens, n_seqs} => {n_embd, n_seq_tokens, n_seqs}
// {d_inner, n_embd} @ {d_inner, n_tokens} => {n_embd, n_tokens}
cur = build_lora_mm(model.layers[il].ssm_out, y, model.layers[il].ssm_out_s);
}
// {n_embd, n_seq_tokens, n_seqs} => {n_embd, n_tokens}
cur = ggml_reshape_2d(ctx0, cur, cur->ne[0], n_seq_tokens * n_seqs);
cb(cur, "mamba_out", il);
return cur;
}
src/models/models.h
+4
View File
@@ -1412,6 +1412,10 @@ struct llama_model_glm4_moe : public llama_model_base {
graph(const llama_model & model, const llm_graph_params & params);
};
struct graph_mtp : public llm_graph_context {
graph_mtp(const llama_model & model, const llm_graph_params & params);
};
std::unique_ptr<llm_graph_context> build_arch_graph(const llm_graph_params & params) const override;
};
tools/parser/template-analysis.cpp → tests/test-chat-analysis.cpp
+11 -7
View File
@@ -11,9 +11,9 @@
#include <vector>
#include <algorithm>
#include "nlohmann/json.hpp"
#include "json.h"
using json = nlohmann::ordered_json;
using json = common_json;
// ANSI color codes - using 256-color palette for brighter colors (all bold)
#define ANSI_RESET "\033[0m"
@@ -84,11 +84,12 @@ static std::string read_file(const std::string & path) {
}
static void print_usage(const char * program_name) {
LOG_ERR("Usage: %s [options]\n", program_name);
LOG_ERR("Debug the auto-parser's differential analysis: render a template with/without tools, reasoning, etc. and show the diffs.\n");
LOG_ERR("\nUsage: %s [options]\n", program_name);
LOG_ERR("\nOptions:\n");
LOG_ERR(" --template <name> Analyze specific template from test suite (e.g., 'deepseek' or 'DeepSeek-V3.1')\n");
LOG_ERR(" --template-file <path> Analyze custom template file\n");
LOG_ERR(" --all Analyze all templates from test suite\n");
LOG_ERR(" --all Analyze all templates from test suite (default when no arguments are given)\n");
LOG_ERR("\nExamples:\n");
LOG_ERR(" %s --all\n", program_name);
LOG_ERR(" %s --template deepseek\n", program_name);
@@ -97,14 +98,17 @@ static void print_usage(const char * program_name) {
static bool parse_options(int argc, char ** argv, analysis_options & opts) {
if (argc < 2) {
print_usage(argv[0]);
return false;
// default mode: analyze all templates from the test suite
opts.analyze_all = true;
}
for (int i = 1; i < argc; ++i) {
std::string arg = argv[i];
if (arg == "--all") {
if (arg == "-h" || arg == "--help") {
print_usage(argv[0]);
return false;
} else if (arg == "--all") {
opts.analyze_all = true;
} else if (arg == "--template") {
if (i + 1 >= argc) {
tools/cli/cli-context.cpp
+5 -6
View File
@@ -6,8 +6,7 @@
#include "log.h"
#include "console.h"
#define JSON_ASSERT GGML_ASSERT
#include <nlohmann/json.hpp>
#include "json.h"
#include <algorithm>
#include <cctype>
@@ -16,7 +15,7 @@
#include <map>
#include <set>
using json = nlohmann::ordered_json;
using json = common_json;
struct cli_context_impl {
json messages = json::array();
@@ -73,7 +72,7 @@ static std::string format_error_message(const json & err) {
// err is the raw response body of a failed request; it may or may not be JSON
static std::string format_error_message(const std::string & err) {
json parsed = json::parse(err, nullptr, false);
json parsed = json::parse_no_throw(err);
if (!parsed.is_discarded()) {
return format_error_message(parsed);
}
@@ -157,7 +156,7 @@ bool cli_context::init() {
if (!list_and_ask_models()) {
return false;
}
} catch (const json::parse_error & e) {
} catch (const common_json_error & e) {
ui::show_error(e.what());
ui::show_message("This might be caused by an incorrect server-base endpoint URL");
return false;
@@ -364,7 +363,7 @@ bool cli_context::generate_completion(generated_content & content_out, cli_timin
ui::assistant_turn a;
std::string err = client.post_sse("/v1/chat/completions", body.dump(), should_stop, [&](const std::string & payload) {
json chunk = json::parse(payload, nullptr, false);
json chunk = json::parse_no_throw(payload);
if (chunk.is_discarded()) {
return;
}
tools/mtmd/clip-graph.h
+9 -3
View File
@@ -137,9 +137,15 @@ struct clip_graph {
int il,
ggml_tensor * sinks = nullptr) const;
// implementation of the 2D RoPE without adding a new op in ggml
// this is not efficient (use double the memory), but works on all backends
// TODO: there was a more efficient which relies on ggml_view and ggml_rope_ext_inplace, but the rope inplace does not work well with non-contiguous tensors ; we should fix that and revert back to the original implementation in https://github.com/ggml-org/llama.cpp/pull/13065
// implementation of the 2D RoPE using two ggml_rope_ext calls
//
// unlike GGML_ROPE_TYPE_VISION which forces NEOX ordering, this rotates adjacent pairs (normal ordering)
//
// example:
// given a single head with size = 8 --> [00000000]
// dims [0, 4) rotate with pos_a, dims [4, 8) rotate with pos_b --> [aaaabbbb]
// interleave_freq = false --> both halves use the same inv_freq set (like GGML_ROPE_TYPE_VISION)
// interleave_freq = true --> first half uses even inv_freq, second half uses odd inv_freq (used by pixtral)
ggml_tensor * build_rope_2d(
ggml_context * ctx0,
ggml_tensor * cur,
tools/mtmd/clip-model.h
+4 -4
View File
@@ -29,10 +29,10 @@ enum patch_merge_type {
PATCH_MERGE_SPATIAL_UNPAD,
};
// all algos are Pillow-compatible (matching PIL.Image.resize output)
enum resize_algo {
RESIZE_ALGO_BILINEAR, // stretch to target resolution
RESIZE_ALGO_BICUBIC, // center-crop when aspect ratio doesn't match
RESIZE_ALGO_BICUBIC_PILLOW,
RESIZE_ALGO_BILINEAR,
RESIZE_ALGO_BICUBIC,
RESIZE_ALGO_LANCZOS,
};
@@ -73,7 +73,7 @@ struct clip_hparams {
int32_t preproc_max_tiles = 0;
int32_t preproc_tile_size = 0; // local tile size (deepseek-ocr)
resize_algo image_resize_algo_rf = RESIZE_ALGO_BICUBIC;
resize_algo image_resize_algo_ov = RESIZE_ALGO_BILINEAR;
resize_algo image_resize_algo_ov = RESIZE_ALGO_BICUBIC;
pad_style image_pad_rf = PAD_CEIL; // padding style for the refined image (e.g. llava-1.6)
pad_style image_pad_ov = PAD_NONE; // padding style for the overview image (e.g. llava-1.6)
std::array<uint8_t, 3> image_pad_color_rf = {0, 0, 0}; // padding color for refined image
tools/mtmd/clip.cpp
+43 -62
View File
@@ -877,8 +877,6 @@ ggml_tensor * clip_graph::build_attn(
}
// implementation of the 2D RoPE without adding a new op in ggml
// this is not efficient (use double the memory), but works on all backends
// TODO: there was a more efficient which relies on ggml_view and ggml_rope_ext_inplace, but the rope inplace does not work well with non-contiguous tensors ; we should fix that and revert back to the original implementation in https://github.com/ggml-org/llama.cpp/pull/13065
ggml_tensor * clip_graph::build_rope_2d(
ggml_context * ctx0,
ggml_tensor * cur,
@@ -887,9 +885,7 @@ ggml_tensor * clip_graph::build_rope_2d(
const float freq_base,
const bool interleave_freq
) {
const int64_t n_dim = cur->ne[0];
const int64_t n_head = cur->ne[1];
const int64_t n_pos = cur->ne[2];
const int64_t n_dim = cur->ne[0];
// for example, if we have cur tensor of shape (n_dim=8, n_head, n_pos)
// we will have a list of 4 inv_freq: 1e-0, 1e-1, 1e-2, 1e-3
@@ -903,46 +899,30 @@ ggml_tensor * clip_graph::build_rope_2d(
? std::pow(freq_base, (float)-2/n_dim)
: 1.0;
// first half
ggml_tensor * first;
{
first = ggml_view_3d(ctx0, cur,
n_dim/2, n_head, n_pos,
cur->nb[1],
cur->nb[2],
0);
first = ggml_rope_ext(
ctx0,
first,
pos_a, // positions
nullptr, // freq factors
n_dim/2, // n_dims
0, 0, freq_base,
1.0f, 0.0f, 1.0f, 0.0f, 0.0f
);
}
// first half, dims [0, n_dim/2)
cur = ggml_rope_ext(
ctx0,
cur,
pos_a, // positions
nullptr, // freq factors
n_dim/2, // n_dims
0, 0, freq_base,
1.0f, 0.0f, 1.0f, 0.0f, 0.0f
);
// second half
ggml_tensor * second;
{
second = ggml_view_3d(ctx0, cur,
n_dim/2, n_head, n_pos,
cur->nb[1],
cur->nb[2],
n_dim/2 * ggml_element_size(cur));
second = ggml_rope_ext(
ctx0,
second,
pos_b, // positions
nullptr, // freq factors
n_dim/2, // n_dims
0, 0, freq_base,
freq_scale_odd,
0.0f, 1.0f, 0.0f, 0.0f
);
}
// second half, dims [n_dim/2, n_dim)
cur = ggml_rope_ext(
ctx0,
cur,
pos_b, // positions
nullptr, // freq factors
n_dim/2, // n_dims
0, 0, freq_base,
freq_scale_odd,
0.0f, 1.0f, 0.0f, 0.0f
);
cur = ggml_rope_set_offset(cur, n_dim/2);
cur = ggml_concat(ctx0, first, second, 0);
return cur;
}
@@ -1521,20 +1501,18 @@ struct clip_model_loader {
hparams.image_pad_color = {122, 116, 104};
if (!hparams.image_res_candidates.empty()) {
hparams.image_resize_pad = PAD_CEIL;
hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
} else {
// llava-1.6 default params
hparams.image_pad_ov = PAD_NONE;
hparams.image_pad_rf = PAD_CEIL;
hparams.image_pad_color_rf = {122, 116, 104};
hparams.image_resize_algo_rf = RESIZE_ALGO_BICUBIC;
hparams.image_resize_algo_ov = RESIZE_ALGO_BILINEAR;
}
} break;
case PROJECTOR_TYPE_GLM_EDGE:
{
hparams.image_resize_pad = PAD_CEIL;
hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
} break;
case PROJECTOR_TYPE_MINICPMV:
{
@@ -1591,6 +1569,7 @@ struct clip_model_loader {
case PROJECTOR_TYPE_IDEFICS3:
{
// use default llava-uhd preprocessing params
hparams.image_resize_algo = RESIZE_ALGO_LANCZOS;
get_u32(KEY_PROJ_SCALE_FACTOR, hparams.n_merge, false);
get_u32(KEY_PREPROC_IMAGE_SIZE, hparams.image_longest_edge, false);
hparams.set_limit_image_tokens();
@@ -1617,7 +1596,7 @@ struct clip_model_loader {
// ref: https://huggingface.co/mistral-community/pixtral-12b/blob/main/preprocessor_config.json
// TODO: verify the image_min_tokens
hparams.n_merge = 1; // the original pixtral does not use patch merging
hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
hparams.rope_theta = 10000.0f;
get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge, false);
hparams.set_limit_image_tokens(8, 1024);
@@ -1645,7 +1624,7 @@ struct clip_model_loader {
case PROJECTOR_TYPE_DOTS3NOTE_V:
{
hparams.rope_theta = 10000.0f;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge);
get_u32(KEY_IMAGE_MIN_PIXELS, hparams.image_min_pixels);
get_u32(KEY_IMAGE_MAX_PIXELS, hparams.image_max_pixels);
@@ -1663,7 +1642,7 @@ struct clip_model_loader {
} break;
case PROJECTOR_TYPE_KIMIVL:
{
hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
hparams.rope_theta = 10000.0f;
get_u32(KEY_PROJ_SCALE_FACTOR, hparams.n_merge, false);
// TODO: check kimivl preprocessor for exact values
@@ -1702,7 +1681,7 @@ struct clip_model_loader {
{
hparams.rope_theta = 100.0f;
hparams.n_merge = 3; // pooling_kernel_size
hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
get_u32(KEY_PROJ_SCALE_FACTOR, hparams.n_merge, false);
if (model.proj_type == PROJECTOR_TYPE_GEMMA4UV) {
// for "unified" variant, we directly use a bigger patch size, because the "token merging" is done directly on conv layer
@@ -1719,6 +1698,7 @@ struct clip_model_loader {
// Gemma3n uses MobileNetV5 which produces 256 tokens (16x16)
// Similar configuration to Gemma3
hparams.n_merge = 1; // MobileNetV5 handles resizing internally
hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
get_u32(KEY_PROJ_SCALE_FACTOR, hparams.n_merge, false);
} break;
case PROJECTOR_TYPE_QWEN2VL:
@@ -1726,7 +1706,7 @@ struct clip_model_loader {
case PROJECTOR_TYPE_QWEN3VL:
{
hparams.n_merge = 2; // default value for Qwen 2 and 2.5
hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge, false);
get_u32(KEY_WIN_ATTN_PATTERN, hparams.n_wa_pattern, model.proj_type == PROJECTOR_TYPE_QWEN25VL); // only 2.5 requires it
// ref: https://huggingface.co/Qwen/Qwen2.5-VL-7B-Instruct/blob/main/preprocessor_config.json
@@ -1747,7 +1727,7 @@ struct clip_model_loader {
case PROJECTOR_TYPE_MINIMAX_M3:
{
hparams.n_merge = 2; // spatial_merge_size
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
hparams.image_resize_pad = PAD_NONE;
get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge, false);
// n_merge is used as a divisor in clip_image_batch_encode
@@ -1772,7 +1752,7 @@ struct clip_model_loader {
case PROJECTOR_TYPE_MIMOVL:
{
hparams.n_merge = 2; // spatial_merge_size
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge, false);
get_u32(string_format(KEY_N_HEAD_KV, "vision"), hparams.n_head_kv);
// 1D banded sliding-window radius (visual_token_window_size); required
@@ -1819,15 +1799,15 @@ struct clip_model_loader {
log_ffn_op = "gelu_erf";
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
// reka model performs better when using resize_bicubic, which stretches
// the image to fit fixed square size
// reka model performs better when the image is stretched to fit
// fixed square size (no padding)
hparams.image_resize_pad = PAD_NONE;
} break;
case PROJECTOR_TYPE_GLM4V:
{
hparams.rope_theta = 10000.0f;
hparams.n_merge = 2; // default value for GLM4-V
hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge, false);
hparams.set_limit_image_tokens(8, 4096);
hparams.set_warmup_n_tokens(46*46); // avoid OOM on warmup
@@ -1835,6 +1815,7 @@ struct clip_model_loader {
case PROJECTOR_TYPE_LLAMA4:
{
hparams.rope_theta = 10000.0f;
hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
get_u32(KEY_PROJ_SCALE_FACTOR, hparams.n_merge, false);
set_llava_uhd_res_candidates(model, 3);
} break;
@@ -1946,7 +1927,7 @@ struct clip_model_loader {
case PROJECTOR_TYPE_PADDLEOCR:
{
hparams.n_merge = 2;
hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
get_u32(KEY_IMAGE_MIN_PIXELS, hparams.image_min_pixels);
get_u32(KEY_IMAGE_MAX_PIXELS, hparams.image_max_pixels);
@@ -1958,7 +1939,7 @@ struct clip_model_loader {
hparams.patch_size = 16;
hparams.image_size = 1024;
hparams.warmup_image_size = 1024;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
hparams.image_pad_color = {127, 127, 127};
get_u32(KEY_SAM_N_BLOCK, hparams.sam_n_layer, true);
@@ -1988,7 +1969,7 @@ struct clip_model_loader {
case PROJECTOR_TYPE_HUNYUANVL:
{
hparams.n_merge = 2;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
hparams.image_resize_algo = RESIZE_ALGO_LANCZOS;
hparams.image_resize_pad = PAD_NONE;
hparams.ffn_op = FFN_GELU;
hparams.set_limit_image_tokens(256, 16384);
@@ -2061,12 +2042,12 @@ struct clip_model_loader {
case PROJECTOR_TYPE_JANUS_PRO:
{
hparams.image_pad_color = {127, 127, 127};
hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
} break;
case PROJECTOR_TYPE_GRANITE4_VISION:
{
// SigLIP tower.
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
hparams.image_resize_pad = PAD_CEIL;
// NOTE: feature_layers loaded in common path as optional
tools/mtmd/models/gemma4v.cpp
+22 -42
View File
@@ -44,51 +44,31 @@ ggml_cgraph * clip_graph_gemma4v::build() {
// similar to build_rope_2d, but use neox ordering
auto add_pos = [&](ggml_tensor * cur, const clip_layer &) {
const int64_t n_dim = cur->ne[0];
const int64_t n_head = cur->ne[1];
const int64_t n_pos = cur->ne[2];
const int64_t n_dim = cur->ne[0];
// first half
ggml_tensor * first;
{
first = ggml_view_4d(ctx0, cur,
n_dim/2, n_head, n_pos, n_batch,
cur->nb[1],
cur->nb[2],
cur->nb[3],
0);
first = ggml_rope_ext(
ctx0,
first,
pos_x, // positions
nullptr, // freq factors
n_dim/2, // n_dims
GGML_ROPE_TYPE_NEOX, 0, hparams.rope_theta,
1.0f, 0.0f, 1.0f, 0.0f, 0.0f
);
}
// first half, dims [0, n_dim/2)
cur = ggml_rope_ext(
ctx0,
cur,
pos_x, // positions
nullptr, // freq factors
n_dim/2, // n_dims
GGML_ROPE_TYPE_NEOX, 0, hparams.rope_theta,
1.0f, 0.0f, 1.0f, 0.0f, 0.0f
);
// second half
ggml_tensor * second;
{
second = ggml_view_4d(ctx0, cur,
n_dim/2, n_head, n_pos, n_batch,
cur->nb[1],
cur->nb[2],
cur->nb[3],
n_dim/2 * ggml_element_size(cur));
second = ggml_rope_ext(
ctx0,
second,
pos_y, // positions
nullptr, // freq factors
n_dim/2, // n_dims
GGML_ROPE_TYPE_NEOX, 0, hparams.rope_theta,
1.0f, 0.0f, 1.0f, 0.0f, 0.0f
);
}
// second half, dims [n_dim/2, n_dim)
cur = ggml_rope_ext(
ctx0,
cur,
pos_y, // positions
nullptr, // freq factors
n_dim/2, // n_dims
GGML_ROPE_TYPE_NEOX, 0, hparams.rope_theta,
1.0f, 0.0f, 1.0f, 0.0f, 0.0f
);
cur = ggml_rope_set_offset(cur, n_dim/2);
cur = ggml_concat(ctx0, first, second, 0);
return cur;
};
tools/mtmd/models/minimax-m3.cpp
+8 -16
View File
@@ -2,30 +2,22 @@
ggml_tensor * clip_graph_minimax_m3::apply_rope(
ggml_tensor * x, ggml_tensor * pos_h, ggml_tensor * pos_w) {
const int64_t Hn = x->ne[1];
const int64_t P = x->ne[2];
const size_t es = ggml_element_size(x);
const int dh = (int) x->ne[0];
const int axd = 2 * ((2 * (dh / 2) / 3) / 2);
const int dh = (int) x->ne[0];
const int axd = 2 * ((2 * (dh / 2) / 3) / 2);
GGML_ASSERT(x->nb[0] == es);
GGML_ASSERT(3 * axd <= dh);
const float th = hparams.rope_theta;
// layout of x is [t, h, w, pad]
// t is unrotated, h and w are rotated, pad is unrotated
// note: everything from n_dims onward untouched, so w and pad are rotated in one call.
auto sl = [&](int off, int n) {
return ggml_cont(ctx0, ggml_view_3d(ctx0, x, n, Hn, P, x->nb[1], x->nb[2], (size_t) off * es));
};
ggml_tensor * t = sl(0, axd);
ggml_tensor * h = sl(axd, axd);
ggml_tensor * w = sl(2 * axd, dh - 2 * axd); // w + pad
x = ggml_rope_ext(ctx0, x, pos_h, nullptr, axd, GGML_ROPE_TYPE_NEOX, 0, th, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f);
x = ggml_rope_set_offset(x, axd);
h = ggml_rope_ext(ctx0, h, pos_h, nullptr, axd, GGML_ROPE_TYPE_NEOX, 0, th, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f);
w = ggml_rope_ext(ctx0, w, pos_w, nullptr, axd, GGML_ROPE_TYPE_NEOX, 0, th, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f);
return ggml_concat(ctx0, ggml_concat(ctx0, t, h, 0), w, 0);
x = ggml_rope_ext(ctx0, x, pos_w, nullptr, axd, GGML_ROPE_TYPE_NEOX, 0, th, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f);
x = ggml_rope_set_offset(x, 2 * axd);
return x;
}
ggml_cgraph * clip_graph_minimax_m3::build() {
tools/mtmd/mtmd-helper.cpp
+42 -9
View File
@@ -43,6 +43,11 @@
#ifdef MTMD_VIDEO
#include "sheredom/subprocess.h"
#include <thread>
#ifndef _WIN32
#include <csignal>
#include <fcntl.h>
#include <pthread.h>
#endif
#endif
//
@@ -523,7 +528,8 @@ struct mtmd_helper_video {
// RAII wrapper for managing subprocess
struct subprocess_handle {
struct subprocess_s proc = {};
bool alive = false;
bool created = false; // process exists and must be cleaned up
bool alive = false; // process can still give us data
std::thread feeder;
subprocess_handle() = default;
@@ -532,18 +538,27 @@ struct mtmd_helper_video {
~subprocess_handle() { stop(); }
void stop() {
if (alive) {
subprocess_terminate(&proc);
// note: alive becomes false on stdout EOF, but the process still needs cleanup
if (!created) {
return;
}
subprocess_terminate(&proc);
#ifdef _WIN32
// no SIGPIPE on windows: a blocked feeder only gets a broken pipe once we close our read end of the child stdin
if (proc.hStdInput) {
CloseHandle(proc.hStdInput);
proc.hStdInput = nullptr;
}
#endif
// join before destroy: feeder holds a FILE* from subprocess_stdin;
// subprocess_destroy closes it, so the thread must finish first
if (feeder.joinable()) {
feeder.join();
}
if (alive) {
subprocess_destroy(&proc);
alive = false;
}
subprocess_join(&proc, nullptr); // reap the child, or else it stays a zombie
subprocess_destroy(&proc);
created = false;
alive = false;
}
FILE * stdout_pipe() {
@@ -553,10 +568,21 @@ struct mtmd_helper_video {
// buf is tied to lifetime of mtmd_helper_video, so it's guaranteed to outlive the feeder thread
void start_feeder(const std::vector<uint8_t> & buf) {
feeder = std::thread([this, &buf]() {
#ifndef _WIN32
// ffmpeg can exit before it reads all the input, for example when ffprobe already got the metadata.
// the write below must then fail with EPIPE, instead of killing the process with SIGPIPE
sigset_t sigpipe_set;
sigemptyset(&sigpipe_set);
sigaddset(&sigpipe_set, SIGPIPE);
pthread_sigmask(SIG_BLOCK, &sigpipe_set, nullptr); // linux sends the signal to the writing thread
#endif
FILE * f = subprocess_stdin(&proc);
if (!f) {
return;
}
#ifdef F_SETNOSIGPIPE
fcntl(fileno(f), F_SETNOSIGPIPE, 1); // macos/bsd send it to the process, so turn it off per fd
#endif
fwrite(buf.data(), 1, buf.size(), f);
fclose(f);
proc.stdin_file = nullptr; // prevent double-close in subprocess_destroy
@@ -602,7 +628,8 @@ struct mtmd_helper_video {
LOG_ERR("%s: failed to launch ffprobe\n", __func__);
return false;
}
probe_sp.alive = true;
probe_sp.created = true;
probe_sp.alive = true;
if (is_buf_input()) {
probe_sp.start_feeder(input_buf);
@@ -674,6 +701,11 @@ struct mtmd_helper_video {
}
cmd.push_back("-nostdin");
if (is_buf_input()) {
// remove the 64KB read-ahead limit of cache:, or else ffmpeg cannot reach a moov atom at end of file
cmd.push_back("-read_ahead_limit");
cmd.push_back("-1");
}
cmd.push_back("-i");
// cache:pipe:0 wraps stdin with a seekable in-memory cache, letting ffmpeg seek
// backwards for container headers (e.g. MP4 moov atom at end of file)
@@ -712,7 +744,8 @@ struct mtmd_helper_video {
subprocess_option_search_user_path | subprocess_option_inherit_environment,
&sp.proc);
sp.alive = (ret == 0);
sp.created = (ret == 0);
sp.alive = (ret == 0);
LOG_DBG("%s: subprocess_create ret=%d proc_alive=%d\n", __func__, ret, (int)sp.alive);
if (sp.alive && is_buf_input()) {
tools/mtmd/mtmd-image.cpp
+82 -244
View File
@@ -58,22 +58,7 @@ struct img_tool {
if (padding == PAD_NONE) {
// direct resize
switch (algo) {
case RESIZE_ALGO_BILINEAR:
resize_bilinear(src, dst, target_resolution.width, target_resolution.height);
break;
case RESIZE_ALGO_BICUBIC:
resize_bicubic(src, dst, target_resolution.width, target_resolution.height);
break;
case RESIZE_ALGO_BICUBIC_PILLOW:
resize_bicubic_pillow(src, dst, target_resolution.width, target_resolution.height);
break;
case RESIZE_ALGO_LANCZOS:
resize_lanczos_pillow(src, dst, target_resolution.width, target_resolution.height);
break;
default:
throw std::runtime_error("Unsupported resize algorithm");
}
resize_pillow(src, dst, target_resolution.width, target_resolution.height, algo);
} else {
// resize with padding
clip_image_u8 resized_image;
@@ -90,22 +75,7 @@ struct img_tool {
new_height = std::min(static_cast<int>(std::ceil(src.get_size().height * scale)), target_resolution.height);
}
switch (algo) {
case RESIZE_ALGO_BILINEAR:
resize_bilinear(src, resized_image, new_width, new_height);
break;
case RESIZE_ALGO_BICUBIC:
resize_bicubic(src, resized_image, new_width, new_height);
break;
case RESIZE_ALGO_BICUBIC_PILLOW:
resize_bicubic_pillow(src, resized_image, new_width, new_height);
break;
case RESIZE_ALGO_LANCZOS:
resize_lanczos_pillow(src, resized_image, new_width, new_height);
break;
default:
throw std::runtime_error("Unsupported resize algorithm");
}
resize_pillow(src, resized_image, new_width, new_height, algo);
// fill dst with pad_color
fill(dst, pad_color);
@@ -224,152 +194,37 @@ struct img_tool {
}
private:
// Bilinear resize function
static void resize_bilinear(const clip_image_u8 & src, clip_image_u8 & dst, int target_width, int target_height) {
const auto src_size = src.get_size();
if (src_size.width == 0 || src_size.height == 0) { dst.set_size({0, 0}, false); return; }
if (target_width <= 0) target_width = 1;
if (target_height <= 0) target_height = 1;
dst.set_size({target_width, target_height}, false);
if (src.is_placeholder()) {
// no-op for placeholder image, just set the size and return
return;
}
float x_ratio = target_width > 1 ? static_cast<float>(src_size.width - 1) / (target_width - 1) : 0.0f;
float y_ratio = target_height > 1 ? static_cast<float>(src_size.height - 1) / (target_height - 1) : 0.0f;
for (int y = 0; y < target_height; ++y) {
for (int x = 0; x < target_width; ++x) {
float px = x * x_ratio;
float py = y * y_ratio;
int x0 = std::min(static_cast<int>(px), src_size.width - 1);
int y0 = std::min(static_cast<int>(py), src_size.height - 1);
int x1 = std::min(x0 + 1, src_size.width - 1);
int y1 = std::min(y0 + 1, src_size.height - 1);
float xf = px - x0;
float yf = py - y0;
const auto p00 = src.get_pixel(x0, y0);
const auto p10 = src.get_pixel(x1, y0);
const auto p01 = src.get_pixel(x0, y1);
const auto p11 = src.get_pixel(x1, y1);
std::array<uint8_t, 3> pixel;
for (int c = 0; c < 3; ++c) {
float top = lerp(static_cast<float>(p00[c]), static_cast<float>(p10[c]), xf);
float bottom = lerp(static_cast<float>(p01[c]), static_cast<float>(p11[c]), xf);
pixel[c] = static_cast<uint8_t>(lerp(top, bottom, yf));
}
dst.set_pixel(x, y, pixel);
}
}
}
// Bicubic resize function
// part of image will be cropped if the aspect ratio is different
static void resize_bicubic(const clip_image_u8 & img, clip_image_u8 & dst, int target_width, int target_height) {
const auto img_size = img.get_size();
const int nx = img_size.width;
const int ny = img_size.height;
dst.set_size({target_width, target_height}, false);
if (img.is_placeholder()) {
// no-op for placeholder image, just set the size and return
return;
}
float Cc;
float C[5] = {};
float d0, d2, d3, a0, a1, a2, a3;
int i, j, k, jj;
int x, y;
float dx, dy;
float tx, ty;
tx = (float)nx / (float)target_width;
ty = (float)ny / (float)target_height;
// Bicubic interpolation; adapted from ViT.cpp, inspired from :
// -> https://github.com/yglukhov/bicubic-interpolation-image-processing/blob/master/libimage.c#L36
// -> https://en.wikipedia.org/wiki/Bicubic_interpolation
for (i = 0; i < target_height; i++) {
for (j = 0; j < target_width; j++) {
x = (int)(tx * j);
y = (int)(ty * i);
dx = tx * j - x;
dy = ty * i - y;
std::array<uint8_t, 3> pixel;
for (k = 0; k < 3; k++) {
for (jj = 0; jj <= 3; jj++) {
d0 = img.get_pixel(clip(x - 1, 0, nx - 1), clip(y - 1 + jj, 0, ny - 1))[k] - img.get_pixel(clip(x, 0, nx - 1), clip(y - 1 + jj, 0, ny - 1))[k];
d2 = img.get_pixel(clip(x + 1, 0, nx - 1), clip(y - 1 + jj, 0, ny - 1))[k] - img.get_pixel(clip(x, 0, nx - 1), clip(y - 1 + jj, 0, ny - 1))[k];
d3 = img.get_pixel(clip(x + 2, 0, nx - 1), clip(y - 1 + jj, 0, ny - 1))[k] - img.get_pixel(clip(x, 0, nx - 1), clip(y - 1 + jj, 0, ny - 1))[k];
a0 = img.get_pixel(clip(x, 0, nx - 1), clip(y - 1 + jj, 0, ny - 1))[k];
a1 = -1.0 / 3 * d0 + d2 - 1.0 / 6 * d3;
a2 = 1.0 / 2 * d0 + 1.0 / 2 * d2;
a3 = -1.0 / 6 * d0 - 1.0 / 2 * d2 + 1.0 / 6 * d3;
C[jj] = a0 + a1 * dx + a2 * dx * dx + a3 * dx * dx * dx;
d0 = C[0] - C[1];
d2 = C[2] - C[1];
d3 = C[3] - C[1];
a0 = C[1];
a1 = -1.0 / 3 * d0 + d2 - 1.0 / 6 * d3;
a2 = 1.0 / 2 * d0 + 1.0 / 2 * d2;
a3 = -1.0 / 6 * d0 - 1.0 / 2 * d2 + 1.0 / 6 * d3;
Cc = a0 + a1 * dy + a2 * dy * dy + a3 * dy * dy * dy;
const uint8_t Cc2 = std::min(std::max(std::round(Cc), 0.0f), 255.0f);
pixel[k] = Cc2;
}
}
dst.set_pixel(j, i, pixel);
}
}
}
// Pillow-compatible separable resampling (Bicubic and Lanczos)
// Pillow-compatible separable resampling (Bilinear, Bicubic and Lanczos)
// Adapted from https://github.com/python-pillow/Pillow/blob/main/src/libImaging/Resample.c
//
// Key properties:
// 1. Separable filtering: horizontal pass followed by vertical pass
// 2. Pre-computes normalized filter coefficients for each output pixel
// 3. Fixed-point integer arithmetic (22 fractional bits) for speed and determinism
static bool resize_bicubic_pillow(const clip_image_u8 & img, clip_image_u8 & dst, int target_width, int target_height) {
return resize_pillow(img, dst, target_width, target_height, /*use_lanczos=*/false);
}
// Lanczos-3 (support radius 3), matches Pillow's Image.LANCZOS
static bool resize_lanczos_pillow(const clip_image_u8 & img, clip_image_u8 & dst, int target_width, int target_height) {
return resize_pillow(img, dst, target_width, target_height, /*use_lanczos=*/true);
}
static bool resize_pillow(
const clip_image_u8 & img,
clip_image_u8 & dst,
int target_width,
int target_height,
bool use_lanczos) {
resize_algo algo) {
// Fixed-point precision: 22 bits = 32 (int32_t) - 8 (uint8_t pixels) - 2 (headroom for accumulation)
// This allows encoding fractional weights as integers: weight * 2^22
const int PRECISION_BITS = 32 - 8 - 2;
// Resample filter: Lanczos-3 (support [-3, 3]) or bicubic with a = -0.5 (support [-2, 2])
// Note: GGML/PyTorch bicubic uses a = -0.75, Pillow uses a = -0.5
// Filter support radius
double filter_support;
switch (algo) {
case RESIZE_ALGO_BILINEAR: filter_support = 1.0; break;
case RESIZE_ALGO_BICUBIC: filter_support = 2.0; break;
case RESIZE_ALGO_LANCZOS: filter_support = 3.0; break;
default:
throw std::runtime_error("Unsupported resize algorithm");
}
// Returns filter weight for distance x from pixel center
auto resample_filter = [use_lanczos](double x) -> double {
if (use_lanczos) {
// Note: for bicubic, Pillow uses a = -0.5 while GGML/PyTorch use a = -0.75
auto resample_filter = [algo](double x) -> double {
if (algo == RESIZE_ALGO_LANCZOS) {
if (-3.0 <= x && x < 3.0) {
auto sinc = [](double v) {
if (v == 0.0) {
@@ -383,10 +238,15 @@ private:
return 0.0;
}
constexpr double a = -0.5;
if (x < 0.0) {
x = -x;
}
if (algo == RESIZE_ALGO_BILINEAR) {
return x < 1.0 ? 1.0 - x : 0.0;
}
constexpr double a = -0.5;
if (x < 1.0) {
return ((a + 2.0) * x - (a + 3.0)) * x * x + 1;
}
@@ -396,9 +256,6 @@ private:
return 0.0; // Zero outside [-2, 2]
};
// Filter support radius: 2 for bicubic, 3 for lanczos
const double filter_support = use_lanczos ? 3.0 : 2.0;
// Clipping function for 8-bit values
auto clip8 = [](int val) -> uint8_t {
if (val < 0) return 0;
@@ -493,100 +350,92 @@ private:
const double fxp_scale = std::ldexp(1.0, PRECISION_BITS); // 1.0 * 2^PRECISION_BITS
for (int i = 0; i < outSize * ksize; i++) {
if (use_lanczos) {
// Pillow adds +/- 0.5 then truncates toward zero; std::round would round twice
const double rounded = pre_weights[i] * fxp_scale + (pre_weights[i] < 0 ? -0.5 : 0.5);
weights[i] = static_cast<int32_t>(rounded);
continue;
}
double tmp_val = pre_weights[i] * fxp_scale;
if (pre_weights[i] < 0) {
tmp_val -= 0.5;
} else {
tmp_val += 0.5;
}
tmp_val = std::round(tmp_val);
tmp_val = std::clamp(tmp_val,
static_cast<double>(std::numeric_limits<int32_t>::min()),
static_cast<double>(std::numeric_limits<int32_t>::max()));
weights[i] = static_cast<int32_t>(tmp_val);
// Pillow adds +/- 0.5 then truncates toward zero; std::round would round twice
const double rounded = pre_weights[i] * fxp_scale + (pre_weights[i] < 0 ? -0.5 : 0.5);
weights[i] = static_cast<int32_t>(rounded);
}
return ksize;
};
// Horizontal resampling pass
// Resizes width from imIn to out_nx, preserving height
auto resample_horizontal = [&](const clip_image_u8 & imIn, clip_image_u8 & imOut,
// Resizes width from src to out_nx, preserving height
auto resample_horizontal = [&](const uint8_t * src, int in_nx, int in_ny,
int out_nx,
int ksize, const std::vector<int> & bounds, const std::vector<int32_t> & weights) {
const int in_ny = imIn.get_size().height;
imOut.set_size({out_nx, in_ny}, false);
std::vector<uint8_t> out((size_t) out_nx * in_ny * 3);
// Process each row independently
for (int yy = 0; yy < in_ny; yy++) {
const uint8_t * src_row = src + (size_t) yy * in_nx * 3;
uint8_t * dst_row = out.data() + (size_t) yy * out_nx * 3;
// For each output pixel in this row
for (int xx = 0; xx < out_nx; xx++) {
// Get the range of input pixels and filter coefficients
int xmin = bounds[xx * 2 + 0]; // First input pixel index
int xcnt = bounds[xx * 2 + 1]; // Number of input pixels
const int xmin = bounds[xx * 2 + 0]; // First input pixel index
const int xcnt = bounds[xx * 2 + 1]; // Number of input pixels
const int32_t * k = &weights[xx * ksize];
const uint8_t * p = src_row + (size_t) xmin * 3;
// Initialize accumulators for RGB channels with rounding bias (0.5 in fixed-point)
// Accumulators for RGB channels, with rounding bias (0.5 in fixed-point)
int32_t ss0 = 1 << (PRECISION_BITS - 1);
int32_t ss1 = 1 << (PRECISION_BITS - 1);
int32_t ss2 = 1 << (PRECISION_BITS - 1);
// Convolve: sum weighted input pixels
for (int x = 0; x < xcnt; x++) {
const auto src_px = imIn.get_pixel(x + xmin, yy);
ss0 += src_px[0] * weights[xx * ksize + x]; // R channel
ss1 += src_px[1] * weights[xx * ksize + x]; // G channel
ss2 += src_px[2] * weights[xx * ksize + x]; // B channel
ss0 += p[0] * k[x];
ss1 += p[1] * k[x];
ss2 += p[2] * k[x];
p += 3;
}
// Convert back from fixed-point (divide by 2^PRECISION_BITS) and clamp to [0,255]
imOut.set_pixel(xx, yy, {clip8(ss0 >> PRECISION_BITS),
clip8(ss1 >> PRECISION_BITS),
clip8(ss2 >> PRECISION_BITS)});
dst_row[xx * 3 + 0] = clip8(ss0 >> PRECISION_BITS);
dst_row[xx * 3 + 1] = clip8(ss1 >> PRECISION_BITS);
dst_row[xx * 3 + 2] = clip8(ss2 >> PRECISION_BITS);
}
}
return out;
};
// Vertical resampling pass
// Resizes height from imIn to out_ny, preserving width
auto resample_vertical = [&](const clip_image_u8 & imIn, clip_image_u8 & imOut,
// Resizes height from src to out_ny, preserving width
// Accumulates whole rows at once (contiguous access, auto-vectorizes well)
auto resample_vertical = [&](const uint8_t * src, int in_nx,
int out_ny,
int ksize, const std::vector<int> & bounds, const std::vector<int32_t> & weight) {
const int in_nx = imIn.get_size().width;
imOut.set_size({in_nx, out_ny}, false);
const size_t row_elems = (size_t) in_nx * 3;
std::vector<uint8_t> out(row_elems * out_ny);
std::vector<int32_t> acc(row_elems);
// For each output row
for (int yy = 0; yy < out_ny; yy++) {
// Get the range of input rows and filter coefficients
int ymin = bounds[yy * 2 + 0]; // First input row index
int ycnt = bounds[yy * 2 + 1]; // Number of input rows
const int ymin = bounds[yy * 2 + 0]; // First input row index
const int ycnt = bounds[yy * 2 + 1]; // Number of input rows
const int32_t * k = &weight[yy * ksize];
// Process each column in this output row
for (int xx = 0; xx < in_nx; xx++) {
// Initialize accumulators for RGB channels with rounding bias
int32_t ss0 = 1 << (PRECISION_BITS - 1);
int32_t ss1 = 1 << (PRECISION_BITS - 1);
int32_t ss2 = 1 << (PRECISION_BITS - 1);
// Rounding bias (0.5 in fixed-point)
std::fill(acc.begin(), acc.end(), 1 << (PRECISION_BITS - 1));
// Convolve: sum weighted input pixels vertically
for (int y = 0; y < ycnt; y++) {
const auto src_px = imIn.get_pixel(xx, y + ymin);
ss0 += src_px[0] * weight[yy * ksize + y]; // R channel
ss1 += src_px[1] * weight[yy * ksize + y]; // G channel
ss2 += src_px[2] * weight[yy * ksize + y]; // B channel
// Convolve: accumulate each weighted input row
for (int y = 0; y < ycnt; y++) {
const uint8_t * src_row = src + (size_t) (ymin + y) * row_elems;
const int32_t w = k[y];
for (size_t i = 0; i < row_elems; i++) {
acc[i] += src_row[i] * w;
}
}
// Convert back from fixed-point and clamp to [0,255]
imOut.set_pixel(xx, yy, {clip8(ss0 >> PRECISION_BITS),
clip8(ss1 >> PRECISION_BITS),
clip8(ss2 >> PRECISION_BITS)});
// Convert back from fixed-point and clamp to [0,255]
uint8_t * dst_row = out.data() + (size_t) yy * row_elems;
for (size_t i = 0; i < row_elems; i++) {
dst_row[i] = clip8(acc[i] >> PRECISION_BITS);
}
}
return out;
};
// Main resampling logic using separable two-pass approach
@@ -610,36 +459,25 @@ private:
}
// Perform two-pass resampling
const uint8_t * src = img.get_ro_buf().data();
if (need_horizontal && need_vertical) {
// Both horizontal and vertical
clip_image_u8 temp;
resample_horizontal(img, temp, target_width, ksize_horiz, bounds_horiz, weights_horiz);
resample_vertical(temp, dst, target_height, ksize_vert, bounds_vert, weights_vert);
auto temp = resample_horizontal(src, src_width, src_height, target_width, ksize_horiz, bounds_horiz, weights_horiz);
dst.set_size({target_width, target_height}, false);
dst.cpy_buf(resample_vertical(temp.data(), target_width, target_height, ksize_vert, bounds_vert, weights_vert));
} else if (need_horizontal) {
// Only horizontal
resample_horizontal(img, dst, target_width, ksize_horiz, bounds_horiz, weights_horiz);
dst.set_size({target_width, src_height}, false);
dst.cpy_buf(resample_horizontal(src, src_width, src_height, target_width, ksize_horiz, bounds_horiz, weights_horiz));
} else if (need_vertical) {
// Only vertical
resample_vertical(img, dst, target_height, ksize_vert, bounds_vert, weights_vert);
dst.set_size({src_width, target_height}, false);
dst.cpy_buf(resample_vertical(src, src_width, target_height, ksize_vert, bounds_vert, weights_vert));
} else {
// No resizing needed - direct copy
dst.set_size(img.get_size(), img.is_placeholder());
if (!img.is_placeholder()) {
dst.cpy_buf(img.get_ro_buf());
}
dst.set_size(img.get_size(), false);
dst.cpy_buf(img.get_ro_buf());
}
return true;
}
static inline int clip(int x, int lower, int upper) {
return std::max(lower, std::min(x, upper));
}
// Linear interpolation between two points
static inline float lerp(float s, float e, float t) {
return s + (e - s) * t;
}
};
@@ -1264,7 +1102,7 @@ mtmd_image_preproc_out mtmd_image_preprocessor_deepseekocr::preprocess(const cli
clip_image_u8 padded;
img_tool::resize(img, padded,
{ base_size, base_size },
RESIZE_ALGO_BICUBIC_PILLOW,
RESIZE_ALGO_BICUBIC,
PAD_NEAREST,
hparams.image_pad_color);
output.append_overview(hparams, padded, true);
@@ -1280,7 +1118,7 @@ mtmd_image_preproc_out mtmd_image_preprocessor_deepseekocr::preprocess(const cli
grid_h = grid.height;
clip_image_u8 refined;
img_tool::resize(img, refined, { tile_size * grid_w, tile_size * grid_h }, RESIZE_ALGO_BICUBIC_PILLOW,
img_tool::resize(img, refined, { tile_size * grid_w, tile_size * grid_h }, RESIZE_ALGO_BICUBIC,
PAD_NONE);
for (int row = 0; row < grid_h; row++) {
tools/parser/CMakeLists.txt
-20
View File
@@ -1,20 +0,0 @@
if (NOT WIN32 OR NOT BUILD_SHARED_LIBS)
# this tool is disabled on Windows when building with shared libraries because it uses internal functions not exported with LLAMA_API
set(TARGET llama-debug-template-parser)
add_executable(${TARGET} debug-template-parser.cpp)
target_link_libraries(${TARGET} PRIVATE llama-common llama ${CMAKE_THREAD_LIBS_INIT})
target_compile_features(${TARGET} PRIVATE cxx_std_17)
if(LLAMA_TOOLS_INSTALL)
install(TARGETS ${TARGET} RUNTIME)
endif()
endif()
set(TARGET llama-template-analysis)
add_executable(${TARGET} template-analysis.cpp)
target_link_libraries(${TARGET} PRIVATE llama-common llama ${CMAKE_THREAD_LIBS_INIT})
target_compile_features(${TARGET} PRIVATE cxx_std_17)
if(LLAMA_TOOLS_INSTALL)
install(TARGETS ${TARGET} RUNTIME)
endif()
tools/parser/debug-template-parser.cpp
-469
View File
@@ -1,469 +0,0 @@
#include "../src/llama-grammar.h"
#include "chat-auto-parser.h"
#include "chat.h"
#include "common.h"
#include "gguf.h"
#include "jinja/runtime.h"
#include "log.h"
#include "nlohmann/json.hpp"
#include "peg-parser.h"
#include <fstream>
#include <iterator>
#include <numeric>
#include <optional>
#include <sstream>
#include <string>
using json = nlohmann::ordered_json;
enum class output_mode {
ANALYSIS, // Only output analysis results (default)
TEMPLATE, // Only output rendered template
BOTH // Output both
};
enum class input_message_type {
NONE, // Don't render any message scenarios (only analysis)
CONTENT_ONLY, // Simple assistant message with content
REASONING_CONTENT, // Message with reasoning_content + content
TOOL_CALL_ONLY, // Message with tool_calls only
CONTENT_TOOL_CALL, // Message with content + tool_calls
REASONING_TOOL_CALL, // Message with reasoning_content + tool_calls
CONTENT_FAKE_TOOL_CALL, // Message with content but no actual tool_calls (for testing)
ALL // Render all scenarios
};
struct debug_options {
std::string template_path;
bool with_tools = true;
bool generation_prompt = true;
bool enable_reasoning = true;
bool debug_jinja = false;
bool force_tool_call = false;
bool parallel_tool_calls = true;
output_mode mode = output_mode::BOTH;
input_message_type input_message = input_message_type::NONE;
};
static std::string read_file(const std::string & path) {
std::ifstream fin(path, std::ios::binary);
if (!fin.is_open()) {
throw std::runtime_error("Could not open file: " + path);
}
std::ostringstream buf;
buf << fin.rdbuf();
return buf.str();
}
static std::string read_gguf_chat_template(const std::string & path) {
struct gguf_init_params params = { /*no_alloc =*/true, // We only need metadata, not tensor data
/*ctx=*/nullptr };
struct gguf_context * ctx = gguf_init_from_file(path.c_str(), params);
if (ctx == nullptr) {
throw std::runtime_error("Could not open GGUF file: " + path);
}
const char * key = "tokenizer.chat_template";
int64_t key_id = gguf_find_key(ctx, key);
if (key_id == -1) {
gguf_free(ctx);
throw std::runtime_error("GGUF file does not contain chat template key: " + std::string(key));
}
const char * template_str = gguf_get_val_str(ctx, key_id);
if (template_str == nullptr) {
gguf_free(ctx);
throw std::runtime_error("GGUF file contains chat template key but value is null");
}
std::string result = template_str;
gguf_free(ctx);
return result;
}
static void print_usage(const char * program_name) {
LOG_ERR("Usage: %s <template_or_gguf_path> [options]\n", program_name);
LOG_ERR("\nOptions:\n");
LOG_ERR(" --no-tools Disable tool definitions\n");
LOG_ERR(" --force-tool-call Set tool calls to forced\n");
LOG_ERR(" --parallel-tool-calls=0|1 Set parallel_tool_calls (default: 1)\n");
LOG_ERR(" --generation-prompt=0|1 Set add_generation_prompt (default: 1)\n");
LOG_ERR(" --enable-reasoning=0|1 Enable reasoning parsing (default: 1)\n");
LOG_ERR(" --output=MODE Output mode: analysis, template, both (default: both)\n");
LOG_ERR(" --debug-jinja Enable Jinja fine-grained debug\n");
LOG_ERR(" --input-message=TYPE Message type to render:\n");
LOG_ERR(" content_only, reasoning_content, tool_call_only,\n");
LOG_ERR(" content_tool_call, reasoning_tool_call,\n");
LOG_ERR(" content_fake_tool_call, all\n");
LOG_ERR("\nExamples:\n");
LOG_ERR(" %s template.jinja --input-message=all --generation-prompt=1\n", program_name);
LOG_ERR(" %s template.jinja --output=template --input-message=tool_call_only\n", program_name);
}
static bool parse_bool_option(const std::string & value) {
return value == "1" || value == "true" || value == "yes";
}
static bool parse_options(int argc, char ** argv, debug_options & opts) {
if (argc < 2) {
print_usage(argv[0]);
return false;
}
opts.template_path = argv[1];
for (int i = 2; i < argc; ++i) {
std::string arg = argv[i];
if (arg == "--force-tool-call") {
opts.force_tool_call = true;
} else if (arg == "--debug-jinja") {
opts.debug_jinja = true;
} else if (arg == "--no-tools") {
opts.with_tools = false;
} else if (arg.rfind("--parallel-tool-calls=", 0) == 0) {
opts.parallel_tool_calls = parse_bool_option(arg.substr(22));
} else if (arg.rfind("--generation-prompt=", 0) == 0) {
opts.generation_prompt = parse_bool_option(arg.substr(20));
} else if (arg.rfind("--enable-reasoning=", 0) == 0) {
opts.enable_reasoning = parse_bool_option(arg.substr(19));
} else if (arg.rfind("--output=", 0) == 0) {
std::string mode = arg.substr(9);
if (mode == "analysis") {
opts.mode = output_mode::ANALYSIS;
} else if (mode == "template") {
opts.mode = output_mode::TEMPLATE;
} else if (mode == "both") {
opts.mode = output_mode::BOTH;
} else {
LOG_ERR("Unknown output mode: %s\n", mode.c_str());
return false;
}
} else if (arg.rfind("--input-message=", 0) == 0) {
std::string type = arg.substr(16);
if (type == "content_only") {
opts.input_message = input_message_type::CONTENT_ONLY;
} else if (type == "reasoning_content") {
opts.input_message = input_message_type::REASONING_CONTENT;
} else if (type == "tool_call_only") {
opts.input_message = input_message_type::TOOL_CALL_ONLY;
} else if (type == "content_tool_call") {
opts.input_message = input_message_type::CONTENT_TOOL_CALL;
} else if (type == "reasoning_tool_call") {
opts.input_message = input_message_type::REASONING_TOOL_CALL;
} else if (type == "content_fake_tool_call") {
opts.input_message = input_message_type::CONTENT_FAKE_TOOL_CALL;
} else if (type == "all") {
opts.input_message = input_message_type::ALL;
} else {
LOG_ERR("Unknown input message type: %s\n", type.c_str());
return false;
}
} else {
LOG_ERR("Unknown option: %s\n", arg.c_str());
print_usage(argv[0]);
return false;
}
}
return true;
}
static json build_user_message() {
return json{
{ "role", "user" },
{ "content", "Hello, please help me with a task." }
};
}
static json build_content_only_message() {
return json{
{ "role", "assistant" },
{ "content", "Hello! I'm here to help you with your task." }
};
}
static json build_reasoning_content_message() {
return json{
{ "role", "assistant" },
{ "content", "Hello! I'm here to help you with your task." },
{ "reasoning_content", "The user is greeting me and asking for help. I should respond politely." }
};
}
static json build_tool_call_only_message() {
return json{
{ "role", "assistant" },
{ "content", nullptr },
{ "tool_calls",
json::array({ json{
{ "type", "function" },
{ "function", json{ { "name", "test_function_name" },
{ "arguments", json::object({ { "param1", "value1" }, { "param2", "value2" } }) } } },
{ "id", "123456789" } } }) }
};
}
static json build_content_tool_call_message() {
return json{
{ "role", "assistant" },
{ "content", "I'll help you by calling a function." },
{ "tool_calls",
json::array({ json{
{ "type", "function" },
{ "function",
json{ { "name", "test_function_name" },
{ "arguments", json::object({ { "param1", "value1" }, { "param2", "value2" } }) } } } } }) }
};
}
static json build_reasoning_tool_call_message() {
return json{
{ "role", "assistant" },
{ "content", nullptr },
{ "reasoning_content", "I need to call a function to help with this task." },
{ "tool_calls",
json::array({ json{
{ "type", "function" },
{ "function",
json{ { "name", "test_function_name" },
{ "arguments", json::object({ { "param1", "value1" }, { "param2", "value2" } }) } } } } }) }
};
}
static json build_content_fake_tool_call_message() {
// This message has content but NO tool_calls field
// It's used to test if a template renders tool definitions but not tool calls
return json{
{ "role", "assistant" },
{ "content", "I'll help you by calling a function." }
};
}
static json build_tools_definition() {
json parameters_schema = json::object();
parameters_schema["type"] = "object";
parameters_schema["properties"] = json::object();
parameters_schema["properties"]["param1"] = json::object({
{ "type", "string" },
{ "description", "First parameter" }
});
parameters_schema["properties"]["param2"] = json::object({
{ "type", "string" },
{ "description", "Second parameter" }
});
parameters_schema["required"] = json::array({ "param1" });
return json::array({
json{ { "type", "function" },
{ "function", json{ { "name", "test_function_name" },
{ "description", "A test function for debugging" },
{ "parameters", parameters_schema } } } }
});
}
static void render_scenario(const common_chat_template & tmpl,
const std::string & scenario_name,
const json & messages,
const json & tools,
bool add_generation_prompt,
bool enable_thinking) {
LOG_ERR("\n=== Scenario: %s ===\n", scenario_name.c_str());
LOG_ERR("add_generation_prompt: %s, enable_thinking: %s\n", add_generation_prompt ? "true" : "false",
enable_thinking ? "true" : "false");
// When add_generation_prompt is true, add a trailing user message to trigger the prompt
json final_messages = messages;
if (add_generation_prompt && !messages.empty() && messages.back().value("role", "") == "assistant") {
final_messages.push_back(json{
{ "role", "user" },
{ "content", "Now please continue with another response." }
});
}
LOG_ERR("Messages:\n%s\n", final_messages.dump(2).c_str());
try {
autoparser::generation_params inputs;
inputs.messages = final_messages;
inputs.add_generation_prompt = add_generation_prompt;
inputs.extra_context["enable_thinking"] = enable_thinking;
if (!tools.is_null() && tools.is_array() && !tools.empty()) {
inputs.tools = tools;
}
std::string output = common_chat_template_direct_apply(tmpl, inputs);
LOG_ERR("\n--- Rendered Output ---\n");
LOG_ERR("%s\n", output.c_str());
LOG_ERR("--- End Output (length: %zu) ---\n", output.length());
} catch (const std::exception & e) {
LOG_ERR("Rendering failed: %s\n", e.what());
}
}
static void render_all_scenarios(const common_chat_template & tmpl,
const json & tools,
bool add_generation_prompt,
bool enable_thinking,
input_message_type message_type) {
json user_msg = build_user_message();
auto render_if = [&](input_message_type type, const std::string & name, const json & assistant_msg) {
if (message_type == input_message_type::ALL || message_type == type) {
json messages = json::array({ user_msg, assistant_msg });
render_scenario(tmpl, name, messages, tools, add_generation_prompt, enable_thinking);
}
};
render_if(input_message_type::CONTENT_ONLY, "content_only", build_content_only_message());
render_if(input_message_type::REASONING_CONTENT, "reasoning_content", build_reasoning_content_message());
render_if(input_message_type::TOOL_CALL_ONLY, "tool_call_only", build_tool_call_only_message());
render_if(input_message_type::CONTENT_TOOL_CALL, "content_tool_call", build_content_tool_call_message());
render_if(input_message_type::REASONING_TOOL_CALL, "reasoning_tool_call", build_reasoning_tool_call_message());
render_if(input_message_type::CONTENT_FAKE_TOOL_CALL, "content_fake_tool_call",
build_content_fake_tool_call_message());
// Also render with add_generation_prompt=true to show the prompt ending
if (message_type == input_message_type::ALL) {
LOG_ERR("\n\n=== Generation Prompt Scenarios (add_generation_prompt=true) ===\n");
json prompt_messages = json::array({ user_msg });
render_scenario(tmpl, "generation_prompt_only", prompt_messages, tools, true, enable_thinking);
// With enable_thinking toggled
render_scenario(tmpl, "generation_prompt_thinking_disabled", prompt_messages, tools, true, false);
}
}
static autoparser::generation_params prepare_params(const debug_options & opts, const json & tools) {
autoparser::generation_params params;
params.messages = json::array({ build_user_message() });
params.reasoning_format = opts.enable_reasoning ? COMMON_REASONING_FORMAT_DEEPSEEK : COMMON_REASONING_FORMAT_NONE;
params.enable_thinking = opts.enable_reasoning;
params.add_generation_prompt = opts.generation_prompt;
if (opts.with_tools) {
params.tools = tools;
params.tool_choice = opts.force_tool_call ? COMMON_CHAT_TOOL_CHOICE_REQUIRED : COMMON_CHAT_TOOL_CHOICE_AUTO;
} else {
params.tools = json();
params.tool_choice = COMMON_CHAT_TOOL_CHOICE_NONE;
}
params.parallel_tool_calls = opts.parallel_tool_calls;
return params;
}
int main(int argc, char ** argv) {
// Set log level to most verbose to capture all debug output
common_log_set_verbosity_thold(99);
debug_options opts;
if (!parse_options(argc, argv, opts)) {
return 1;
}
if (opts.debug_jinja || std::getenv("LLAMA_DEBUG_JINJA") != nullptr) {
jinja::enable_debug(true);
}
std::string template_source;
try {
// Check if the file is a GGUF file
if (opts.template_path.size() >= 5 &&
opts.template_path.compare(opts.template_path.size() - 5, 5, ".gguf") == 0) {
template_source = read_gguf_chat_template(opts.template_path);
} else {
template_source = read_file(opts.template_path);
}
} catch (const std::exception & e) {
LOG_ERR("Error reading template: %s\n", e.what());
return 1;
}
LOG_ERR("Analyzing template: %s\n", opts.template_path.c_str());
LOG_ERR("Options: with_tools=%s, generation_prompt=%s, enable_reasoning=%s\n", opts.with_tools ? "true" : "false",
opts.generation_prompt ? "true" : "false", opts.enable_reasoning ? "true" : "false");
try {
common_chat_template chat_template(template_source, "", "");
json tools = opts.with_tools ? build_tools_definition() : json();
autoparser::generation_params params = prepare_params(opts, tools);
common_chat_params parser_data;
if (std::optional<common_chat_params> spec_tmpl =
common_chat_try_specialized_template(chat_template, template_source, params)) {
LOG_ERR("\n");
LOG_ERR("This template uses a specialized parser, analysis results will not be available.\n");
parser_data = *spec_tmpl;
} else {
// Render template scenarios if requested
if (opts.input_message != input_message_type::NONE &&
(opts.mode == output_mode::TEMPLATE || opts.mode == output_mode::BOTH)) {
LOG_ERR("\n");
LOG_ERR("================================================================================\n");
LOG_ERR(" TEMPLATE RENDERING OUTPUT\n");
LOG_ERR("================================================================================\n");
render_all_scenarios(chat_template, tools, opts.generation_prompt, opts.enable_reasoning,
opts.input_message);
}
// Output analysis if requested
if (opts.mode == output_mode::ANALYSIS || opts.mode == output_mode::BOTH) {
LOG_ERR("\n");
LOG_ERR("================================================================================\n");
LOG_ERR(" TEMPLATE ANALYSIS\n");
LOG_ERR("================================================================================\n");
autoparser::autoparser analysis;
analysis.analyze_template(chat_template);
// Generate Parser
parser_data = autoparser::peg_generator::generate_parser(chat_template, params, analysis);
}
}
if (!std::empty(parser_data.parser)) {
LOG_ERR("\n=== Generated Parser ===\n");
common_peg_arena arena;
arena.load(parser_data.parser);
LOG_ERR("%s\n", arena.dump(arena.root()).c_str());
LOG_ERR("\n=== Generated Grammar ===\n");
LOG_ERR("%s\n", parser_data.grammar.c_str());
LOG_ERR("\n=== Generated Lazy Grammar ===\n");
LOG_ERR("%d\n", parser_data.grammar_lazy);
LOG_ERR("\n=== Generated Grammar Triggers ===\n");
for (const common_grammar_trigger & cgt : parser_data.grammar_triggers) {
LOG_ERR("Token: %d | Type: %d | Value: %s\n", cgt.token, cgt.type, cgt.value.c_str());
}
LOG_ERR("\n=== Preserved Tokens ===\n");
for (const std::string & token : parser_data.preserved_tokens) {
LOG_ERR(" '%s'\n", token.c_str());
}
if (!parser_data.grammar.empty()) {
LOG_ERR("\n=== Verifying created grammar ===\n");
auto * grammar = llama_grammar_init_impl(nullptr, parser_data.grammar.c_str(), "root",
parser_data.grammar_lazy, nullptr, 0, nullptr, 0);
if (grammar != nullptr) {
LOG_ERR("\n=== Grammar successfully created ===\n");
}
}
}
} catch (const std::exception & e) {
LOG_ERR("Analysis failed: %s\n", e.what());
return 1;
}
return 0;
}
tools/server/server-chat.cpp
+1 -1
View File
@@ -153,7 +153,7 @@ json server_chat_convert_responses_to_chatcmpl(const json & response_body) {
prev_msg["content"] = json::array();
}
auto & prev_content = prev_msg["content"];
prev_content.insert(prev_content.end(), chatcmpl_content.begin(), chatcmpl_content.end());
prev_content.insert(chatcmpl_content);
} else {
item.erase("status");
item.erase("type");
tools/server/server-chat.h
+1 -3
View File
@@ -6,9 +6,7 @@
#include "server-common.h"
#include "server-http.h"
#include <nlohmann/json_fwd.hpp>
using json = nlohmann::ordered_json;
#include "json.h"
// Convert OpenAI Responses API format to OpenAI Chat Completions API format
json server_chat_convert_responses_to_chatcmpl(const json & body);
tools/server/server-common.cpp
+1 -1
View File
@@ -1540,7 +1540,7 @@ std::vector<llama_token_data> get_token_probabilities(llama_context * ctx, int i
}
std::string safe_json_to_str(const json & data) {
return data.dump(-1, ' ', false, json::error_handler_t::replace);
return data.dump_safe();
}
// TODO: reuse llama_detokenize
tools/server/server-common.h
+5 -6
View File
@@ -6,8 +6,7 @@
#include "chat.h"
#include "mtmd.h"
#define JSON_ASSERT GGML_ASSERT
#include <nlohmann/json.hpp>
#include "json.h"
#include <atomic>
#include <chrono>
@@ -19,7 +18,7 @@
#include <string>
#include <vector>
using json = nlohmann::ordered_json;
using json = common_json;
#define SLT_DBG(slot, fmt, ...) LOG_DBG("slot %12.*s: id %2d | task %d | " fmt, 12, __func__, (slot).id, ((slot).task ? (slot).task->id : -1), __VA_ARGS__)
#define SLT_TRC(slot, fmt, ...) LOG_TRC("slot %12.*s: id %2d | task %d | " fmt, 12, __func__, (slot).id, ((slot).task ? (slot).task->id : -1), __VA_ARGS__)
@@ -42,9 +41,9 @@ static T json_value(const json & body, const std::string & key, const T & defaul
// Fallback null to default value
if (body.contains(key) && !body.at(key).is_null()) {
try {
return body.at(key);
} catch (NLOHMANN_JSON_NAMESPACE::detail::type_error const & err) {
LOG_WRN("Wrong type supplied for parameter '%s'. Expected '%s', using default value: %s\n", key.c_str(), json(default_value).type_name(), err.what());
return body.at(key).get<T>();
} catch (const common_json_error & err) {
LOG_WRN("Wrong type supplied for parameter '%s', using default value: %s\n", key.c_str(), err.what());
return default_value;
}
} else {
tools/server/server-context.cpp
+29 -19
View File
@@ -35,8 +35,6 @@
#include <windows.h>
#endif
using json = nlohmann::ordered_json;
constexpr int HTTP_POLLING_SECONDS = 1;
static common_speculative_output_limits server_output_limits(const common_params & params) {
@@ -657,14 +655,14 @@ struct server_slot {
res["n_prompt_tokens_processed"] = stats.n_prompt_processed;
res["n_prompt_tokens_cache"] = stats.n_prompt_cached;
res["params"] = ptask->params.to_json(only_metrics);
res["next_token"] = {
res["next_token"] = json::array({
{
{"has_next_token", has_next_token},
{"has_new_line", has_new_line},
{"n_remain", n_remaining()},
{"n_decoded", stats.n_gen},
}
};
});
if (!only_metrics) {
res["prompt"] = ptask->tokens.detokenize(ctx_tgt, true);
@@ -860,8 +858,10 @@ private:
// slots / clients
std::vector<server_slot> slots;
int trace = 0;
int slots_debug = 0;
int trace = 0; // env: LLAMA_TRACE
int slots_debug = 0; // env: LLAMA_SERVER_SLOTS_DEBUG
int slots_n_diff = 0; // env: LLAMA_SERVER_SLOTS_N_DIFF
int n_empty_consecutive = 0;
std::unique_ptr<server_prompt_cache> prompt_cache;
@@ -1249,6 +1249,15 @@ private:
}
}
{
const char * LLAMA_SERVER_SLOTS_N_DIFF = getenv("LLAMA_SERVER_SLOTS_N_DIFF");
slots_n_diff = LLAMA_SERVER_SLOTS_N_DIFF ? atoi(LLAMA_SERVER_SLOTS_N_DIFF) : 0;
if (slots_n_diff) {
SRV_WRN("LLAMA_SERVER_SLOTS_N_DIFF = %d\n", slots_n_diff);
}
}
// the update_slots() logic will always submit a maximum of n_batch or n_parallel tokens
// note that n_batch can be > n_ctx (e.g. for non-causal attention models such as BERT where the KV cache is not used)
{
@@ -3181,8 +3190,8 @@ private:
// when the prompt prefix does not match, print the tokens around the mismatch
// this is useful for debugging prompt caching
if (slots_debug) {
const int np0 = std::max<int>(n_past - 4, 0);
const int np1 = std::min<int>(n_past + 6, std::min(slot.prompt.tokens.size(), slot.task->tokens.size()));
const int np0 = std::max<int>(n_past - slots_n_diff, 0);
const int np1 = std::min<int>(n_past + slots_n_diff + 2, std::min(slot.prompt.tokens.size(), slot.task->tokens.size()));
std::stringstream ss0;
std::stringstream ss1;
@@ -4165,7 +4174,8 @@ std::unique_ptr<server_res_generator> server_routes::handle_completions_impl(
// tasks.reserve(inputs.size()); // TODO: this is inaccurate due to child tasks
// message delimiters for checkpointing
auto delimiters = common_chat_msg_delimiters_parse(json_value(data, "message_delimiters", json::array()));
json delims = json_value(data, "message_delimiters", json::array());
auto delimiters = common_chat_msg_delimiters_parse(delims);
delimiters.tokenize(ctx_server.vocab);
for (size_t i = 0; i < inputs.size(); i++) {
@@ -4428,8 +4438,8 @@ static json get_res_model_info(const server_context_meta & meta) {
static json get_res_models(const server_context_meta & meta) {
// note: do NOT use ctx_server here, otherwise it's not possible to use this during sleep
return {
{"models", {
return json{
{"models", json::array({
{
{"name", meta.model_name},
{"model", meta.model_name},
@@ -4438,23 +4448,23 @@ static json get_res_models(const server_context_meta & meta) {
{"digest", ""}, // dummy value, llama.cpp does not support managing model file's hash
{"type", "model"},
{"description", ""},
{"tags", {""}},
{"capabilities", meta.has_mtmd ? json({"completion","multimodal"}) : json({"completion"})},
{"tags", json::array({""})},
{"capabilities", meta.has_mtmd ? json::array({"completion","multimodal"}) : json::array({"completion"})},
{"parameters", ""},
{"details", {
{"parent_model", ""},
{"format", "gguf"},
{"family", ""},
{"families", {""}},
{"families", json::array({""})},
{"parameter_size", ""},
{"quantization_level", ""}
}}
}
}},
})},
{"object", "list"},
{"data", {
{"data", json::array({
get_res_model_info(meta),
}}
})}
};
}
@@ -4990,7 +5000,7 @@ void server_routes::init_routes() {
std::string content;
if (body.count("tokens") != 0) {
const llama_tokens tokens = body.at("tokens");
const llama_tokens tokens = body.at("tokens").get<llama_tokens>();
content = tokens_to_str(ctx_server.vocab, tokens);
}
@@ -5297,7 +5307,7 @@ std::unique_ptr<server_res_generator> server_routes::handle_embeddings_impl(cons
int embd_normalize = params.embd_normalize;
if (body.count("embd_normalize") != 0) {
embd_normalize = body.at("embd_normalize");
embd_normalize = body.at("embd_normalize").get<int>();
if (meta->pooling_type == LLAMA_POOLING_TYPE_NONE) {
SRV_DBG("embd_normalize is not supported by pooling type %d, ignoring it\n", meta->pooling_type);
}
tools/server/server-context.h
+1 -1
View File
@@ -4,7 +4,7 @@
#include "server-task.h"
#include "server-queue.h"
#include <nlohmann/json_fwd.hpp>
#include "json.h"
#include <cstddef>
#include <memory>
tools/server/server-models.cpp
+1 -1
View File
@@ -2462,7 +2462,7 @@ server_http_proxy::server_http_proxy(
bool has_files = !files.empty();
if (has_files) {
json form_fields = json::parse(body, nullptr, false);
json form_fields = json::parse_no_throw(body);
if (!form_fields.is_discarded()) {
auto boundary = generate_multipart_boundary();
effective_body = build_multipart_body(form_fields, files, boundary);
tools/server/server-schema.cpp
+2 -3
View File
@@ -503,7 +503,7 @@ std::vector<std::unique_ptr<field>> make_llama_cmpl_schema(const common_params &
->set_handler([&](field_eval_context & ctx, const json & data) {
const auto & samplers = data.at("samplers");
if (samplers.is_array()) {
ctx.params.sampling.samplers = common_sampler_types_from_names(samplers);
ctx.params.sampling.samplers = common_sampler_types_from_names(samplers.get<std::vector<std::string>>());
} else if (samplers.is_string()) {
ctx.params.sampling.samplers = common_sampler_types_from_chars(samplers.get<std::string>());
}
@@ -580,8 +580,7 @@ static void handle_with_catch(const char * name, std::function<void()> func) {
// treat a null value as absent so clients can send null to request the server default
static bool has_value(const json & data, const char * n) {
auto it = data.find(n);
return it != data.end() && !it->is_null();
return data.contains(n) && !data.at(n).is_null();
}
template <typename T>
tools/server/server-task.cpp
+13 -15
View File
@@ -12,8 +12,6 @@
#include <sstream>
using json = nlohmann::ordered_json;
//
// task_params
//
@@ -304,7 +302,7 @@ json completion_token_output::probs_vector_to_json(const std::vector<completion_
}
float completion_token_output::logarithm(float x) {
// nlohmann::json converts -inf to null, so we need to prevent that
// the JSON library converts -inf to null, so we need to prevent that
return x == 0.0f ? std::numeric_limits<float>::lowest() : std::log(x);
}
@@ -407,7 +405,7 @@ json server_task_result_cmpl_final::to_json_oaicompat() {
res["__verbose"] = to_json_non_oaicompat();
}
if (stats.is_set()) {
res.push_back({"timings", stats.to_json()});
res["timings"] = stats.to_json();
}
return res;
@@ -455,7 +453,7 @@ json server_task_result_cmpl_final::to_json_oaicompat_chat() {
res["__verbose"] = to_json_non_oaicompat();
}
if (stats.is_set()) {
res.push_back({"timings", stats.to_json()});
res["timings"] = stats.to_json();
}
return res;
@@ -516,7 +514,7 @@ json server_task_result_cmpl_final::to_json_oaicompat_chat_stream() {
}
if (stats.is_set()) {
deltas.back().push_back({"timings", stats.to_json()});
deltas.back()["timings"] = stats.to_json();
}
// extra fields for debugging purposes
@@ -709,7 +707,7 @@ json server_task_result_cmpl_final::to_json_oaicompat_resp_stream() {
});
if (stats.is_set()) {
server_sent_events.back().at("data").push_back({"timings", stats.to_json()});
server_sent_events.back().at("data")["timings"] = stats.to_json();
}
return server_sent_events;
@@ -1061,10 +1059,10 @@ json server_task_result_cmpl_partial::to_json_non_oaicompat() {
};
// populate the timings object when needed (usually for the last response or with timings_per_token enabled)
if (stats.is_set()) {
res.push_back({"timings", stats.to_json()});
res["timings"] = stats.to_json();
}
if (is_progress) {
res.push_back({"prompt_progress", progress.to_json()});
res["prompt_progress"] = progress.to_json();
}
if (!prob_output.probs.empty()) {
res["completion_probabilities"] = completion_token_output::probs_vector_to_json({prob_output}, post_sampling_probs);
@@ -1101,10 +1099,10 @@ json server_task_result_cmpl_partial::to_json_oaicompat() {
res["__verbose"] = to_json_non_oaicompat();
}
if (stats.is_set()) {
res.push_back({"timings", stats.to_json()});
res["timings"] = stats.to_json();
}
if (is_progress) {
res.push_back({"prompt_progress", progress.to_json()});
res["prompt_progress"] = progress.to_json();
}
return res;
@@ -1155,10 +1153,10 @@ json server_task_result_cmpl_partial::to_json_oaicompat_chat() {
}
if (stats.is_set()) {
last_json.push_back({"timings", stats.to_json()});
last_json["timings"] = stats.to_json();
}
if (is_progress) {
last_json.push_back({"prompt_progress", progress.to_json()});
last_json["prompt_progress"] = progress.to_json();
}
}
@@ -1305,10 +1303,10 @@ json server_task_result_cmpl_partial::to_json_oaicompat_resp() {
if (!events.empty()) {
json & data = events.back().at("data");
if (stats.is_set()) {
data.push_back({"timings", stats.to_json()});
data["timings"] = stats.to_json();
}
if (is_progress) {
data.push_back({"prompt_progress", progress.to_json()});
data["prompt_progress"] = progress.to_json();
}
}
tools/server/server-task.h
-1
View File
@@ -11,7 +11,6 @@
// TODO: prevent including the whole server-common.h as we only use server_tokens
#include "server-common.h"
using json = nlohmann::ordered_json;
enum server_task_type {
SERVER_TASK_TYPE_COMPLETION,
tools/server/server-tools.cpp
+1 -1
View File
@@ -2156,7 +2156,7 @@ void server_tools::setup(const std::vector<std::string> & enabled_tools,
res->status = 200;
res->data = safe_json_to_str(result);
}
} catch (const json::exception & e) {
} catch (const common_json_error & e) {
res->status = 400;
res->data = safe_json_to_str(format_error_response(e.what(), ERROR_TYPE_INVALID_REQUEST));
} catch (const std::invalid_argument & e) {
tools/server/tests/unit/test_slot_save.py
+20 -1
View File
@@ -319,7 +319,6 @@ def test_slot_save_restore_with_two_images(mmproj_server):
"prompt": prompt,
})
assert res.status_code == 200
content = res.body["content"]
prompt_n_full = res.body["timings"]["prompt_n"]
assert prompt_n_full > 64
@@ -345,6 +344,26 @@ def test_slot_save_restore_with_two_images(mmproj_server):
assert res.status_code == 200
assert res.body["timings"]["cache_n"] == prompt_n_full - 1
assert res.body["timings"]["prompt_n"] == 1
content = res.body["content"]
res = server.make_request("POST", "/slots/1?action=restore", data={
"filename": "mm_slot_two_images.bin",
})
assert res.status_code == 200
assert res.body["n_restored"] == n_saved
res = server.make_request("POST", "/completions", data={
"temperature": 0.0,
"top_k": 1,
"id_slot": 0,
"cache_prompt": True,
"prompt": prompt,
})
assert res.status_code == 200
assert res.body["timings"]["cache_n"] == prompt_n_full - 1
assert res.body["timings"]["prompt_n"] == 1
content = res.body["content"]
assert res.body["content"] == content
tools/server/tests/unit/test_vision_api.py
+1 -1
View File
@@ -121,7 +121,7 @@ def test_vision_chat_completion_token_count():
"prompt, image_data, success, re_content",
[
# test model is trained on CIFAR-10, but it's quite dumb due to small size
("What is this: <__media__>\n", "IMG_BASE64_0", True, "(cat)+"),
("What is this: <__media__>\n", "IMG_BASE64_0", True, "(cat)+|(automobile)+"),
("What is this: <__media__>\n", "IMG_BASE64_1", True, "(frog)+"),
("What is this: <__media__>\n", "malformed", False, None), # non-image data
("What is this:\n", "", False, None), # empty string
tools/server/tests/utils.py
+1 -1
View File
@@ -623,7 +623,7 @@ class ServerPreset:
server.model_hf_repo = "ggml-org/tinygemma3-GGUF:Q8_0"
server.model_alias = "tinygemma3"
server.n_ctx = 1024
server.n_batch = 32
server.n_batch = 512
server.n_slots = 2
server.n_predict = 4
server.seed = 42
tools/ui/src/lib/components/app/chat/ChatAttachments/ChatAttachmentsList/ChatAttachmentsList.svelte
+5 -4
View File
@@ -3,8 +3,9 @@
ChatAttachmentsListItem,
DialogChatAttachmentsPreview,
DialogMcpResourcePreview,
HorizontalScrollCarousel
ScrollCarousel
} from '$lib/components/app';
import { ScrollCarouselVariant } from '$lib/enums';
import type { DatabaseMessageExtraMcpResource } from '$lib/types';
import { getAttachmentDisplayItems, isMcpPrompt, isMcpResource } from '$lib/utils';
@@ -42,7 +43,7 @@
uploadedFiles = $bindable([])
}: Props = $props();
let carouselRef: HorizontalScrollCarousel | undefined = $state();
let carouselRef: ScrollCarousel | undefined = $state();
let mcpResourcePreviewOpen = $state(false);
let mcpResourcePreviewExtra = $state<DatabaseMessageExtraMcpResource | null>(null);
let previewFocusIndex = $state(0);
@@ -91,11 +92,11 @@
{#if displayItems.length > 0}
<div class={className} {style}>
{#if limitToSingleRow}
<HorizontalScrollCarousel bind:this={carouselRef}>
<ScrollCarousel bind:this={carouselRef} variant={ScrollCarouselVariant.CENTER}>
{#each displayItems as item (item.id)}
{@render attachmentitem(item)}
{/each}
</HorizontalScrollCarousel>
</ScrollCarousel>
{:else}
<div class="flex flex-wrap items-start justify-end gap-3">
{#each displayItems as item (item.id)}
tools/ui/src/lib/components/app/chat/ChatAttachments/ChatAttachmentsPreview/ChatAttachmentsPreviewThumbnailStrip.svelte
+4 -3
View File
@@ -1,7 +1,8 @@
<script lang="ts">
import { FileText, Music, Video } from '@lucide/svelte';
import { HorizontalScrollCarousel } from '$lib/components/app/misc';
import { ScrollCarousel } from '$lib/components/app';
import { ICON_CLASS_DEFAULT, UI_DATA_ATTRS } from '$lib/constants';
import { ScrollCarouselVariant } from '$lib/enums';
interface PreviewItem {
id: string;
@@ -33,7 +34,7 @@
{#if items.length > 1}
<div class="sticky bottom-0 z-10 mt-4 flex-shrink-0">
<HorizontalScrollCarousel class="max-w-full">
<ScrollCarousel class="max-w-full" variant={ScrollCarouselVariant.CENTER}>
{#each items as item, index (item.id)}
<button
{...{ [UI_DATA_ATTRS.THUMBNAIL_INDEX]: index }}
@@ -64,6 +65,6 @@
{/if}
</button>
{/each}
</HorizontalScrollCarousel>
</ScrollCarousel>
</div>
{/if}
tools/ui/src/lib/components/app/chat/ChatForm/ChatFormInput/ChatFormInputRich.svelte
+1 -1
View File
@@ -625,7 +625,7 @@
}
if (rootElement && (event.key === 'ArrowLeft' || event.key === 'ArrowRight')) {
const isWordJump = (event.altKey || event.ctrlKey) && !event.metaKey;
const isWordJump = (event.altKey || event.ctrlKey) && !event.metaKey && !event.shiftKey;
const isPlainLeft =
event.key === 'ArrowLeft' && !event.altKey && !event.ctrlKey && !event.metaKey;
tools/ui/src/lib/components/app/chat/ChatForm/ChatFormMcpResourcesList.svelte
+4 -6
View File
@@ -1,8 +1,6 @@
<script lang="ts">
import {
ChatAttachmentsListItemMcpResource,
HorizontalScrollCarousel
} from '$lib/components/app';
import { ChatAttachmentsListItemMcpResource, ScrollCarousel } from '$lib/components/app';
import { ScrollCarouselVariant } from '$lib/enums';
import { mcpStore } from '$lib/stores';
interface Props {
@@ -26,7 +24,7 @@
{#if hasAttachments}
<div class={className}>
<HorizontalScrollCarousel gapSize="2">
<ScrollCarousel gapSize="2" variant={ScrollCarouselVariant.CENTER}>
{#each attachments as attachment, i (attachment.id)}
<ChatAttachmentsListItemMcpResource
class={i === 0 ? 'ml-3' : ''}
@@ -35,6 +33,6 @@
onclick={() => handleResourceClick(attachment.resource.uri)}
/>
{/each}
</HorizontalScrollCarousel>
</ScrollCarousel>
</div>
{/if}
tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistant.svelte
+2 -2
View File
@@ -196,7 +196,7 @@
--assistant-min-height-offset: calc(
var(--last-user-message-height, 19rem) + var(--chat-form-height, 6rem) +
var(--chat-form-bottom-position, 0.5rem) + var(--chat-form-padding-top, 6rem) +
var(--assistant-margin-top, 3rem)
var(--assistant-margin-top, 3rem) + var(--chat-tabs-offset, 0px)
);
min-height: calc(100dvh - var(--assistant-min-height-offset));
@@ -204,7 +204,7 @@
--assistant-min-height-offset: calc(
var(--last-user-message-height, 18rem) + var(--chat-form-height, 6rem) +
var(--chat-form-bottom-position, 1rem) + var(--chat-form-padding-top, 6rem) +
var(--assistant-margin-top, 3rem)
var(--assistant-margin-top, 3rem) + var(--chat-tabs-offset, 0px)
);
}
}
tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreen.svelte
+2 -5
View File
@@ -41,10 +41,7 @@
let showDeleteDialog = $state(false);
let showEmptyFileDialog = $state(false);
let isEmpty = $derived(
showCenteredEmpty &&
!conversationsStore.activeConversation &&
conversationsStore.activeMessages.length === 0 &&
!chatStore.isLoading
showCenteredEmpty && conversationsStore.activeMessages.length === 0 && !chatStore.isLoading
);
let activeErrorDialog = $derived(chatStore.errorDialogState);
let isServerLoading = $derived(serverStore.loading);
@@ -297,7 +294,7 @@
<ServerLoadingSplash />
{:else}
<div
class="chat-screen flex grow flex-col min-h-[calc(100dvh-1rem)] md:min-h-full px-4 md:py-0 pt-12 pb-48 md:pb-4"
class="chat-screen flex grow flex-col min-h-[calc(100dvh-1rem)] md:min-h-[calc(100dvh-1rem-var(--chat-tabs-offset,0px))] px-4 md:py-0 pt-12 pb-48 md:pb-4"
style:--chat-form-bottom-position={chatFormBottomPosition}
ondragenter={dragAndDrop.dragHandlers.dragenter}
ondragleave={dragAndDrop.dragHandlers.dragleave}
tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenActionScrollDown.svelte
+1 -1
View File
@@ -14,6 +14,6 @@
tooltip="Scroll to bottom"
size="lg"
iconSize={ICON_CLASS_DEFAULT}
class="h-9 w-9 rounded-full bg-accent text-accent-foreground absolute bottom-4 shadow-md"
class="h-9 w-9 rounded-full bg-muted/60 border border-border/20 shadow-sm text-accent-foreground absolute bottom-4"
/>
</div>
tools/ui/src/lib/components/app/chat/ChatTabs/ChatTabs.svelte
+136
View File
@@ -0,0 +1,136 @@
<script lang="ts">
import ChatTabsItem from './ChatTabsItem.svelte';
import ChatTabsNewChatButton from './ChatTabsNewChatButton.svelte';
import { page } from '$app/state';
import { ScrollCarousel } from '$lib/components/app';
import {
CHAT_TABS_MAX_WIDTH,
NEW_CHAT_LABEL,
NEW_CHAT_TAB_ID,
UI_DATA_ATTRS,
UNNAMED_CHAT_LABEL
} from '$lib/constants';
import { useScrollCarousel } from '$lib/hooks/use-scroll-carousel.svelte';
import { chatStore, conversationsStore, tabsStore, uiStore } from '$lib/stores';
import { tick } from 'svelte';
const carousel = useScrollCarousel();
let activeId = $derived(page.params.id ?? NEW_CHAT_TAB_ID);
let tabs = $derived(
tabsStore.openTabs.map((id) => ({
id,
isNewChat: id === NEW_CHAT_TAB_ID,
name:
id === NEW_CHAT_TAB_ID
? NEW_CHAT_LABEL
: (conversationsStore.conversations.find((c) => c.id === id)?.name ?? UNNAMED_CHAT_LABEL)
}))
);
// hide the New chat button while a new-chat tab is already open
let showNewChatButton = $derived(!tabsStore.openTabs.includes(NEW_CHAT_TAB_ID));
let loadingIds = $derived(new Set(chatStore.getAllLoadingChats()));
function handleClose(id: string) {
void tabsStore.close(id, activeId);
}
function handleStop(id: string, event: MouseEvent) {
event.stopPropagation();
void chatStore.stopGenerationForChat(id);
}
function handleAuxClick(id: string, event: MouseEvent) {
// middle-click closes, like browser tabs
if (event.button === 1) {
event.preventDefault();
handleClose(id);
}
}
let previousTabIds = new Set<string>();
let previousActiveId: string | null = null;
$effect(() => {
const currentIds = new Set(tabs.map((t) => t.id));
const hasAddedTab = tabs.some((t) => !previousTabIds.has(t.id));
previousTabIds = currentIds;
const activeChanged = activeId !== previousActiveId;
previousActiveId = activeId;
// scroll when the active tab changes (a click) or when a new tab is added
if (!hasAddedTab && !activeChanged) return;
// wait for the new tab to be laid out before scrolling to it
void tick().then(() => {
const el = carousel.scrollContainer?.querySelector<HTMLElement>(
`[${UI_DATA_ATTRS.ACTIVE_TAB}]`
);
if (el) {
carousel.scrollToCenter(el);
}
});
});
</script>
<nav
class="group sticky pl-1 top-0 z-10 hidden md:block chat-tabs-fade transition-[padding] duration-200 ease-in-out pt-3.25 {uiStore.isSidebarExpanded
? CHAT_TABS_MAX_WIDTH.EXPANDED_SIDEBAR
: CHAT_TABS_MAX_WIDTH.COLLAPSED_SIDEBAR}"
aria-label="Open conversations"
>
<div class="relative">
<ScrollCarousel
class="h-10"
containerClass="flex h-10 min-w-0 items-center"
innerClass="items-center gap-1.25"
{carousel}
>
{#each tabs as tab (tab.id)}
<ChatTabsItem
{tab}
isActive={tab.id === activeId}
isLoading={loadingIds.has(tab.id)}
onActivate={(id) => tabsStore.activate(id)}
onClose={handleClose}
onStop={handleStop}
onAuxClick={handleAuxClick}
/>
{/each}
{#if showNewChatButton}
<ChatTabsNewChatButton onclick={() => void conversationsStore.openNewChat()} />
{/if}
</ScrollCarousel>
<div
class="pointer-events-none absolute inset-y-0 left-0 z-[5] w-8 bg-gradient-to-r from-background to-transparent transition-opacity {carousel.canScrollLeft
? 'opacity-100'
: 'opacity-0'}"
></div>
<div
class="pointer-events-none absolute inset-y-0 right-0 z-[5] w-8 bg-gradient-to-l from-background to-transparent transition-opacity {carousel.canScrollRight
? 'opacity-100'
: 'opacity-0'}"
></div>
</div>
</nav>
<style>
.chat-tabs-fade {
background: linear-gradient(
to bottom,
color-mix(in srgb, var(--background) 100%, transparent) 25%,
color-mix(in srgb, var(--background) 80%, transparent) 50%,
color-mix(in srgb, var(--background) 40%, transparent) 75%,
transparent 100%
);
}
</style>
tools/ui/src/lib/components/app/chat/ChatTabs/ChatTabsItem.svelte
+156
View File
@@ -0,0 +1,156 @@
<script lang="ts">
import { Loader2, Square, SquarePen, X } from '@lucide/svelte';
import * as Tooltip from '$lib/components/ui/tooltip';
import { cn } from '$lib/components/ui/utils';
import { ICON_CLASS_SM, ICON_CLASS_XS, ROUTES, UI_DATA_ATTRS } from '$lib/constants';
import { RouterService } from '$lib/services/router.service';
interface Tab {
id: string;
isNewChat: boolean;
name: string;
}
interface Props {
tab: Tab;
isActive?: boolean;
isLoading?: boolean;
onActivate?: (id: string) => void;
onClose?: (id: string) => void;
onStop?: (id: string, event: MouseEvent) => void;
onAuxClick?: (id: string, event: MouseEvent) => void;
}
let {
isActive = false,
isLoading = false,
onActivate,
onAuxClick,
onClose,
onStop,
tab
}: Props = $props();
let contentOpacity = $derived(isActive ? '' : 'opacity-45 group-hover:opacity-75');
let href = $derived(tab.isNewChat ? ROUTES.START : RouterService.chat(tab.id));
function handleActivate(event: MouseEvent) {
// let cmd/ctrl/middle-click fall through so the browser keeps its own
// behavior (open in a new window); route the plain click ourselves so the
// new-chat sentinel and history behave exactly like programmatic nav
if (event.metaKey || event.ctrlKey || event.button === 1) return;
event.preventDefault();
onActivate?.(tab.id);
}
// stop/close sit on top of the tab link; swallow their clicks so they do
// not also navigate
function handleActionClick(event: MouseEvent, action: () => void) {
event.preventDefault();
event.stopPropagation();
action();
}
</script>
<!-- the tab link covers the whole item; stop/close sit on top as siblings so
interactive elements are never nested inside the anchor -->
<div
{...{ [UI_DATA_ATTRS.ACTIVE_TAB]: isActive ? 'true' : undefined }}
class={cn(
'relative flex h-8 max-w-52 min-w-0 shrink-0 items-center gap-1 rounded-lg pr-1 text-sm whitespace-nowrap border backdrop-blur-xl first:ml-2',
isLoading ? 'pl-1' : 'pl-3',
isActive
? 'bg-muted/60 border-border/10 shadow-sm text-accent-foreground hover:bg-primary/15'
: 'border-transparent hover:bg-primary/10 hover:border-border/10 hover:shadow-sm'
)}
>
<a
{href}
class="absolute inset-0 z-0 rounded-lg"
onclick={handleActivate}
onauxclick={(e) => onAuxClick?.(tab.id, e)}
aria-current={isActive ? 'page' : undefined}
aria-label={tab.name}
></a>
{#if isLoading}
<Tooltip.Root>
<Tooltip.Trigger>
{#snippet child({ props })}
<button
{...props}
class="stop-button relative z-10 flex h-5 w-5 shrink-0 cursor-pointer items-center justify-center rounded-sm text-muted-foreground transition-colors hover:text-foreground"
onclick={(e) => handleActionClick(e, () => onStop?.(tab.id, e))}
aria-label="Stop generation"
>
<Loader2
class="loading-icon {ICON_CLASS_SM} animate-spin transition-opacity duration-300 {contentOpacity}"
/>
<Square
class="stop-icon hidden {ICON_CLASS_XS} fill-current text-destructive transition-opacity {contentOpacity}"
/>
</button>
{/snippet}
</Tooltip.Trigger>
<Tooltip.Content>
<p>Stop generation</p>
</Tooltip.Content>
</Tooltip.Root>
{/if}
{#if tab.isNewChat}
<SquarePen
class="pointer-events-none {ICON_CLASS_SM} shrink-0 transition-opacity {contentOpacity}"
/>
{/if}
<span class="pointer-events-none truncate transition-opacity {contentOpacity}">{tab.name}</span>
<Tooltip.Root>
<Tooltip.Trigger>
{#snippet child({ props })}
<button
{...props}
class={cn(
'relative z-10 flex h-5 w-5 shrink-0 cursor-pointer items-center justify-center rounded-sm text-muted-foreground transition-opacity hover:bg-foreground/10 hover:text-foreground',
contentOpacity
)}
onclick={(e) => handleActionClick(e, () => onClose?.(tab.id))}
aria-label="Close tab"
>
<X class={ICON_CLASS_SM} />
</button>
{/snippet}
</Tooltip.Trigger>
<Tooltip.Content>
<p>Close tab</p>
</Tooltip.Content>
</Tooltip.Root>
</div>
<style>
.stop-button {
:global(.stop-icon) {
display: none;
}
:global(.loading-icon) {
display: block;
}
&:is(:hover) {
:global(.stop-icon) {
display: block;
}
:global(.loading-icon) {
display: none;
}
}
}
</style>
tools/ui/src/lib/components/app/chat/ChatTabs/ChatTabsNewChatButton.svelte
+30
View File
@@ -0,0 +1,30 @@
<script lang="ts">
import { Plus } from '@lucide/svelte';
import * as Tooltip from '$lib/components/ui/tooltip';
import { ICON_CLASS_DEFAULT } from '$lib/constants';
interface Props {
onclick?: () => void;
}
let { onclick }: Props = $props();
</script>
<Tooltip.Root>
<Tooltip.Trigger>
{#snippet child({ props })}
<button
{...props}
class="backdrop-blur-lg flex h-8 w-8 mr-4 shrink-0 cursor-pointer items-center justify-center rounded-md transition-colors hover:bg-foreground/5"
{onclick}
aria-label="New chat"
>
<Plus class="{ICON_CLASS_DEFAULT} opacity-40 transition-opacity group-hover:opacity-100" />
</button>
{/snippet}
</Tooltip.Trigger>
<Tooltip.Content>
<p>New chat</p>
</Tooltip.Content>
</Tooltip.Root>
tools/ui/src/lib/components/app/chat/index.ts
+12
View File
@@ -686,6 +686,18 @@ export { default as ChatMessageSystem } from './ChatMessages/ChatMessage/ChatMes
*/
export { default as ChatScreen } from './ChatScreen/ChatScreen.svelte';
/**
* **ChatTabs** - Browser-style tab bar for open conversations
*
* Horizontal strip of tabs rendered above ChatScreen in the chat layout,
* one per conversation tracked by tabsStore. The active tab follows the
* route's conversation id; clicking a tab navigates to it, middle-click or
* the close button closes it (switching to the left neighbor when closing
* the active tab), and a trailing "+" button starts a new chat. Shows a
* spinner on tabs with a running generation. Desktop-only.
*/
export { default as ChatTabs } from './ChatTabs/ChatTabs.svelte';
/**
* Visual overlay displayed when user drags files over the chat screen.
* Shows drop zone indicator to guide users where to release files.
tools/ui/src/lib/components/app/misc/HorizontalScrollCarousel.svelte
-96
View File
@@ -1,96 +0,0 @@
<script lang="ts">
import { ChevronLeft, ChevronRight } from '@lucide/svelte';
import { ICON_CLASS_DEFAULT } from '$lib/constants';
import type { Snippet } from 'svelte';
interface Props {
class?: string;
children?: Snippet;
gapSize?: string;
onScrollableChange?: (isScrollable: boolean) => void;
}
let { children, class: className = '', gapSize = '3', onScrollableChange }: Props = $props();
let canScrollLeft = $state(false);
let canScrollRight = $state(false);
let scrollContainer: HTMLDivElement | undefined = $state();
function scrollLeft(event?: MouseEvent) {
event?.stopPropagation();
event?.preventDefault();
if (!scrollContainer) return;
scrollContainer.scrollBy({ behavior: 'smooth', left: scrollContainer.clientWidth * -0.67 });
}
function scrollRight(event?: MouseEvent) {
event?.stopPropagation();
event?.preventDefault();
if (!scrollContainer) return;
scrollContainer.scrollBy({ behavior: 'smooth', left: scrollContainer.clientWidth * 0.67 });
}
function updateScrollButtons() {
if (!scrollContainer) return;
const { clientWidth, scrollLeft, scrollWidth } = scrollContainer;
canScrollLeft = scrollLeft > 0;
canScrollRight = scrollLeft < scrollWidth - clientWidth - 1;
const isScrollable = scrollWidth > clientWidth;
onScrollableChange?.(isScrollable);
}
export function resetScroll() {
if (scrollContainer) {
scrollContainer.scrollLeft = 0;
setTimeout(() => {
updateScrollButtons();
}, 0);
}
}
$effect(() => {
if (!scrollContainer) return;
const observer = new ResizeObserver(() => updateScrollButtons());
observer.observe(scrollContainer);
return () => observer.disconnect();
});
</script>
<div class="relative {className}">
<button
class="absolute top-1/2 left-4 z-10 flex h-6 w-6 -translate-y-1/2 items-center justify-center rounded-full bg-background/25 shadow-md backdrop-blur-xs transition-opacity hover:bg-background/45 disabled:pointer-events-none disabled:opacity-0"
onclick={scrollLeft}
disabled={!canScrollLeft}
aria-label="Scroll left"
>
<ChevronLeft class={ICON_CLASS_DEFAULT} />
</button>
<div
class="scrollbar-hide flex items-start gap-{gapSize} overflow-x-auto"
bind:this={scrollContainer}
onscroll={updateScrollButtons}
>
{@render children?.()}
</div>
<button
class="absolute top-1/2 right-4 z-10 flex h-6 w-6 -translate-y-1/2 items-center justify-center rounded-full bg-background/25 shadow-md backdrop-blur-xs transition-opacity hover:bg-background/45 disabled:pointer-events-none disabled:opacity-0"
onclick={scrollRight}
disabled={!canScrollRight}
aria-label="Scroll right"
>
<ChevronRight class={ICON_CLASS_DEFAULT} />
</button>
</div>
tools/ui/src/lib/components/app/misc/ScrollCarousel.svelte
+131
View File
@@ -0,0 +1,131 @@
<script lang="ts">
import { ChevronLeft, ChevronRight } from '@lucide/svelte';
import { cn } from '$lib/components/ui/utils';
import { ICON_CLASS_DEFAULT } from '$lib/constants';
import { ScrollCarouselVariant } from '$lib/enums';
import { useScrollCarousel } from '$lib/hooks/use-scroll-carousel.svelte';
import type { Snippet } from 'svelte';
interface Props {
children: Snippet;
/** External carousel hook for callers that need to drive it (e.g. scrollToCenter). */
carousel?: ReturnType<typeof useScrollCarousel>;
/** Classes for the outer relative wrapper. */
class?: string;
/** Classes for the scrollable overflow container. */
containerClass?: string;
/** Classes for the min-w-max content wrapper. */
innerClass?: string;
/** Tailwind gap class applied to the content wrapper. */
gapSize?: string;
/** Show the arrows whenever the content overflows, even without hover. */
alwaysShowArrows?: boolean;
/** Arrow placement and styling. */
variant?: ScrollCarouselVariant;
}
let {
alwaysShowArrows = false,
carousel: externalCarousel,
children,
class: className = '',
containerClass = '',
gapSize = '3',
innerClass = '',
variant = ScrollCarouselVariant.TOP
}: Props = $props();
const internalCarousel = useScrollCarousel();
const carousel = $derived(externalCarousel ?? internalCarousel);
const isCenter = $derived(variant === ScrollCarouselVariant.CENTER);
function scrollLeft(event?: MouseEvent) {
event?.stopPropagation();
event?.preventDefault();
const container = carousel.scrollContainer;
if (!container) return;
container.scrollBy({ behavior: 'smooth', left: -(container.clientWidth * 0.67) });
}
function scrollRight(event?: MouseEvent) {
event?.stopPropagation();
event?.preventDefault();
const container = carousel.scrollContainer;
if (!container) return;
container.scrollBy({ behavior: 'smooth', left: container.clientWidth * 0.67 });
}
export function resetScroll() {
const container = carousel.scrollContainer;
if (!container) return;
container.scrollLeft = 0;
setTimeout(() => carousel.updateScrollButtons(), 0);
}
</script>
<div
class={cn('group relative', !isCenter && 'flex items-center', className)}
style={!isCenter ? 'scroll-padding: 1rem;' : undefined}
>
<button
class={cn(
'absolute z-10 flex h-6 w-6 items-center justify-center rounded-full shadow-md transition-opacity',
isCenter
? 'top-1/2 left-4 -translate-y-1/2 bg-background/25 backdrop-blur-xs hover:bg-background/45 disabled:pointer-events-none disabled:opacity-0'
: 'left-2 bg-muted backdrop-blur-sm hover:bg-accent',
!isCenter &&
(carousel.canScrollLeft
? alwaysShowArrows
? 'opacity-100'
: 'opacity-0 group-hover:opacity-100'
: 'pointer-events-none opacity-0')
)}
{...isCenter ? { disabled: !carousel.canScrollLeft } : {}}
onclick={scrollLeft}
aria-label="Scroll left"
>
<ChevronLeft class={ICON_CLASS_DEFAULT} />
</button>
<div
class={cn('scrollbar-hide overflow-x-auto', containerClass)}
bind:this={carousel.scrollContainer}
onscroll={carousel.updateScrollButtons}
>
<div
class={cn('flex min-w-max', isCenter && 'items-start', `gap-${gapSize}`, innerClass)}
bind:this={carousel.contentContainer}
>
{@render children?.()}
</div>
</div>
<button
class={cn(
'absolute z-10 flex h-6 w-6 items-center justify-center rounded-full shadow-md transition-opacity',
isCenter
? 'top-1/2 right-4 -translate-y-1/2 bg-background/25 backdrop-blur-xs hover:bg-background/45 disabled:pointer-events-none disabled:opacity-0'
: 'right-2 bg-muted backdrop-blur-sm hover:bg-accent',
!isCenter &&
(carousel.canScrollRight
? alwaysShowArrows
? 'opacity-100'
: 'opacity-0 group-hover:opacity-100'
: 'pointer-events-none opacity-0')
)}
{...isCenter ? { disabled: !carousel.canScrollRight } : {}}
onclick={scrollRight}
aria-label="Scroll right"
>
<ChevronRight class={ICON_CLASS_DEFAULT} />
</button>
</div>
tools/ui/src/lib/components/app/misc/index.ts
+7 -7
View File
@@ -21,13 +21,6 @@
*/
export { default as ConversationSelection } from './ConversationSelection.svelte';
/**
* Horizontal scrollable carousel with navigation arrows.
* Used for displaying items in a horizontally scrollable container
* with left/right navigation buttons that appear on hover.
*/
export { default as HorizontalScrollCarousel } from './HorizontalScrollCarousel.svelte';
/**
* **TruncatedText** - Text with ellipsis and tooltip
*
@@ -44,6 +37,13 @@ export { default as TruncatedText } from './TruncatedText.svelte';
*/
export { default as KeyboardShortcutInfo } from './KeyboardShortcutInfo.svelte';
/**
* **ScrollCarousel** - Feature/carousel with center-aligned overflow controls
*
* Horizontal scrollable container with arrows that center the focused item.
*/
export { default as ScrollCarousel } from './ScrollCarousel.svelte';
/**
* **CodeBlockActions** - Actions bar for code blocks (copy, preview)
*
tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigation.svelte
+25 -24
View File
@@ -14,7 +14,7 @@
import { useKeyboardShortcuts } from '$lib/hooks/use-keyboard-shortcuts.svelte';
import { useMarqueeSelection } from '$lib/hooks/use-marquee-selection.svelte';
import { RouterService } from '$lib/services/router.service';
import { chatStore, conversationsStore, deviceStore, settingsStore } from '$lib/stores';
import { chatStore, conversationsStore, deviceStore, settingsStore, uiStore } from '$lib/stores';
import { buildConversationTree } from '$lib/utils';
import { circIn } from 'svelte/easing';
import { SvelteSet } from 'svelte/reactivity';
@@ -31,30 +31,29 @@
toggleSidebar: () => toggleExpandedMode()
});
let isExpandedMode = $state(false);
let hoveredTooltip = $state<string | null>(null);
let logoHovered = $state(false);
const isStripExpanded = $derived(isExpandedMode || hoveredTooltip !== null);
const isStripExpanded = $derived(uiStore.isSidebarExpanded || hoveredTooltip !== null);
const isOnMobile = $derived(deviceStore.isMobile);
const alwaysShowOnDesktop = $derived(settingsStore.config.alwaysShowSidebarOnDesktop as boolean);
$effect(() => {
if (alwaysShowOnDesktop && !isOnMobile) {
isExpandedMode = true;
uiStore.isSidebarExpanded = true;
}
});
function toggleExpandedMode() {
isExpandedMode = !isExpandedMode;
uiStore.isSidebarExpanded = !uiStore.isSidebarExpanded;
if (!isExpandedMode) {
if (!uiStore.isSidebarExpanded) {
hoveredTooltip = null;
}
}
$effect(() => {
if (!isExpandedMode) {
if (!uiStore.isSidebarExpanded) {
isSearchModeActive = false;
searchQuery = '';
@@ -66,7 +65,7 @@
$effect(() => {
if (deviceStore.isMobile && page.url.hash.includes(ROUTES.SEARCH)) {
isExpandedMode = false;
uiStore.isSidebarExpanded = false;
}
});
@@ -294,7 +293,7 @@
}
pendingCollapse = setTimeout(() => {
isExpandedMode = false;
uiStore.isSidebarExpanded = false;
pendingCollapse = null;
}, 100);
}
@@ -314,7 +313,7 @@
class={[
'fixed md:sticky top-2 left-2 md:left-0 md:ml-2 md:mt-2 pt-2 z-10 w-[calc(100dvw-1rem)]',
'md:h-[calc(100dvh-1.125rem)]',
isExpandedMode &&
uiStore.isSidebarExpanded &&
(deviceStore.isStandalone
? 'h-[calc(100dvh-2rem)]'
: deviceStore.isIOSDevice
@@ -323,9 +322,9 @@
'rounded-3xl md:rounded-2xl',
'flex flex-col justify-between',
'md:transition-[width,padding] duration-200 ease-out',
isStripExpanded && 'md:w-72 md:bg-muted/60 md:backdrop-blur-xl border-border shadow-md',
isStripExpanded && 'md:w-72 md:bg-muted/60 md:backdrop-blur-xl shadow-md',
!isStripExpanded && 'md:w-12',
isExpandedMode && 'is-expanded'
uiStore.isSidebarExpanded && 'is-expanded'
]}
>
<div class="px-2 flex items-center justify-between">
@@ -337,24 +336,26 @@
onmouseleave={() => (logoHovered = false)}
>
<ActionIcon
icon={!isExpandedMode && logoHovered && innerWidth > 768 ? PanelLeftOpen : Logo}
icon={!uiStore.isSidebarExpanded && logoHovered && innerWidth > 768
? PanelLeftOpen
: Logo}
size="lg"
iconSize="h-4.5 w-4.5 md:h-4 md:w-4"
class="{isExpandedMode
class="{uiStore.isSidebarExpanded
? 'bg-muted! md:bg-foreground/5!'
: 'bg-transparent!'} md:h-9 md:w-9 h-10 w-10 rounded-full md:hover:bg-foreground/10! pointer-events-auto"
href={isExpandedMode ? ROUTES.START : undefined}
onclick={isExpandedMode ? undefined : toggleExpandedMode}
tooltip={isExpandedMode ? undefined : 'Open Sidebar'}
href={uiStore.isSidebarExpanded ? ROUTES.START : undefined}
onclick={uiStore.isSidebarExpanded ? undefined : toggleExpandedMode}
tooltip={uiStore.isSidebarExpanded ? undefined : 'Open Sidebar'}
tooltipSide={TooltipSide.RIGHT}
ariaLabel={isExpandedMode ? 'Go to start' : 'Expand navigation'}
ariaLabel={uiStore.isSidebarExpanded ? 'Go to start' : 'Expand navigation'}
/>
</div>
{#if isOnMobile || (isExpandedMode && !alwaysShowOnDesktop)}
{#if isOnMobile || (uiStore.isSidebarExpanded && !alwaysShowOnDesktop)}
<div
class="flex items-center transition-all duration-150 ease-out {deviceStore.isMobile &&
!isExpandedMode
!uiStore.isSidebarExpanded
? 'opacity-0 h-0!'
: ''}"
in:fade={{ delay: 50, duration: 150, easing: circIn }}
@@ -377,12 +378,12 @@
<div
class="mt-2 flex min-h-0 flex-1 flex-col gap-4 md:gap-1 {deviceStore.isMobile
? 'transition-[opacity,height] duration-200 ease-out'
: ''} {deviceStore.isMobile && !isExpandedMode ? 'opacity-0 !h-0' : ''}"
: ''} {deviceStore.isMobile && !uiStore.isSidebarExpanded ? 'opacity-0 !h-0' : ''}"
in:fade={{ duration: 200 }}
out:fade={{ duration: 200 }}
>
<SidebarNavigationActions
isExpandedMode={innerWidth > 768 ? isExpandedMode : true}
isExpandedMode={innerWidth > 768 ? uiStore.isSidebarExpanded : true}
class="px-2"
bind:isSearchModeActive
bind:searchQuery
@@ -391,7 +392,7 @@
searchQuery = '';
}}
onSearchClick={() => {
isExpandedMode = true;
uiStore.isSidebarExpanded = true;
isSearchModeActive = true;
}}
onNewChat={() => {
@@ -401,7 +402,7 @@
}}
/>
{#if isExpandedMode || isOnMobile}
{#if uiStore.isSidebarExpanded || isOnMobile}
<div class="flex min-h-0 flex-1 flex-col overflow-y-auto">
<SidebarNavigationConversationList
class="px-2"
tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigationActions.svelte
+30 -18
View File
@@ -11,8 +11,8 @@
ROUTES,
SIDEBAR_ACTIONS_ITEMS
} from '$lib/constants';
import { TooltipSide } from '$lib/enums';
import { deviceStore } from '$lib/stores';
import { SidebarAction, TooltipSide } from '$lib/enums';
import { conversationsStore, deviceStore } from '$lib/stores';
import type { Component } from 'svelte';
import { onMount } from 'svelte';
import { circIn } from 'svelte/easing';
@@ -109,14 +109,20 @@
{@const isActive = isItemActive(item)}
{@const isSearchOnMobile = item.icon === Search && deviceStore.isMobile}
{@const itemHref = isSearchOnMobile ? ROUTES.SEARCH : item.route}
{@const itemOnClick = item.route
? () => {
onNewChat?.();
goto(item.route!);
}
: isSearchOnMobile
? undefined
: onSearchClick}
{@const itemOnClick =
item.action === SidebarAction.NEW_CHAT
? () => {
onNewChat?.();
void conversationsStore.openNewChat();
}
: item.route
? () => {
onNewChat?.();
goto(item.route!);
}
: isSearchOnMobile
? undefined
: onSearchClick}
{@const itemTransition = {
delay: !initialized ? i * ICON_STRIP_TRANSITION_DELAY_MULTIPLIER : 0,
duration: ICON_STRIP_TRANSITION_DURATION,
@@ -157,14 +163,20 @@
{#each SIDEBAR_ACTIONS_ITEMS as item, i (item.tooltip)}
{@const isActive = isItemActive(item)}
{@const isSearchOnMobile = item.icon === Search && deviceStore.isMobile}
{@const itemOnClick = item.route
? () => {
onNewChat?.();
goto(item.route!);
}
: isSearchOnMobile
? undefined
: onSearchClick}
{@const itemOnClick =
item.action === SidebarAction.NEW_CHAT
? () => {
onNewChat?.();
void conversationsStore.openNewChat();
}
: item.route
? () => {
onNewChat?.();
goto(item.route!);
}
: isSearchOnMobile
? undefined
: onSearchClick}
{@const itemTransition = {
delay: !initialized ? i * ICON_STRIP_TRANSITION_DELAY_MULTIPLIER : 0,
duration: ICON_STRIP_TRANSITION_DURATION,
tools/ui/src/lib/components/app/settings/SettingsChatMobileHeader.svelte
+39 -66
View File
@@ -1,5 +1,6 @@
<script lang="ts">
import { ChevronLeft, ChevronRight, Settings } from '@lucide/svelte';
import { Settings } from '@lucide/svelte';
import { ScrollCarousel } from '$lib/components/app';
import { ICON_CLASS_DEFAULT, UI_DATA_ATTRS } from '$lib/constants';
import { BooleanString } from '$lib/enums';
import { useScrollCarousel } from '$lib/hooks/use-scroll-carousel.svelte';
@@ -44,70 +45,42 @@
</div>
<div class="border-b border-border/30 py-2">
<div class="relative flex items-center" style="scroll-padding: 1rem;">
<button
class="absolute left-2 z-10 flex h-6 w-6 items-center justify-center rounded-full bg-muted shadow-md backdrop-blur-sm transition-opacity hover:bg-accent {carousel.canScrollLeft
? 'opacity-100'
: 'pointer-events-none opacity-0'}"
onclick={carousel.scrollLeft}
aria-label="Scroll left"
>
<ChevronLeft class={ICON_CLASS_DEFAULT} />
</button>
<div
class="scrollbar-hide overflow-x-auto py-2"
bind:this={carousel.scrollContainer}
onscroll={carousel.updateScrollButtons}
>
<div class="flex min-w-max gap-2">
{#each sections as section (section.title)}
{#if getHref}
<a
class="flex cursor-pointer items-center gap-2 rounded-lg px-3 py-2 text-sm whitespace-nowrap no-underline transition-colors first:ml-4 last:mr-4 hover:bg-accent {isActive(
section
)
? 'bg-accent text-accent-foreground'
: 'text-muted-foreground'}"
{...{ [UI_DATA_ATTRS.ACTIVE]: isActive(section) }}
href={getHref(section)}
onclick={(e: MouseEvent) => {
carousel.scrollToCenter(e.currentTarget as HTMLElement);
}}
>
<section.icon class="{ICON_CLASS_DEFAULT} flex-shrink-0" />
<span>{section.title}</span>
</a>
{:else}
<button
class="flex cursor-pointer items-center gap-2 rounded-lg px-3 py-2 text-sm whitespace-nowrap transition-colors first:ml-4 last:mr-4 hover:bg-accent {isActive(
section
)
? 'bg-accent text-accent-foreground'
: 'text-muted-foreground'}"
{...{ [UI_DATA_ATTRS.ACTIVE]: isActive(section) }}
onclick={(e: MouseEvent) => {
onSectionChange?.(section.title);
carousel.scrollToCenter(e.currentTarget as HTMLElement);
}}
>
<section.icon class="{ICON_CLASS_DEFAULT} flex-shrink-0" />
<span>{section.title}</span>
</button>
{/if}
{/each}
</div>
</div>
<button
class="absolute right-2 z-10 flex h-6 w-6 items-center justify-center rounded-full bg-muted shadow-md backdrop-blur-sm transition-opacity hover:bg-accent {carousel.canScrollRight
? 'opacity-100'
: 'pointer-events-none opacity-0'}"
onclick={carousel.scrollRight}
aria-label="Scroll right"
>
<ChevronRight class={ICON_CLASS_DEFAULT} />
</button>
</div>
<ScrollCarousel {carousel} alwaysShowArrows containerClass="py-2" innerClass="gap-2">
{#each sections as section (section.title)}
{#if getHref}
<a
class="flex cursor-pointer items-center gap-2 rounded-lg px-3 py-2 text-sm whitespace-nowrap no-underline transition-colors first:ml-4 last:mr-4 hover:bg-accent {isActive(
section
)
? 'bg-accent text-accent-foreground'
: 'text-muted-foreground'}"
{...{ [UI_DATA_ATTRS.ACTIVE]: isActive(section) }}
href={getHref(section)}
onclick={(e: MouseEvent) => {
carousel.scrollToCenter(e.currentTarget as HTMLElement);
}}
>
<section.icon class="{ICON_CLASS_DEFAULT} flex-shrink-0" />
<span>{section.title}</span>
</a>
{:else}
<button
class="flex cursor-pointer items-center gap-2 rounded-lg px-3 py-2 text-sm whitespace-nowrap transition-colors first:ml-4 last:mr-4 hover:bg-accent {isActive(
section
)
? 'bg-accent text-accent-foreground'
: 'text-muted-foreground'}"
{...{ [UI_DATA_ATTRS.ACTIVE]: isActive(section) }}
onclick={(e: MouseEvent) => {
onSectionChange?.(section.title);
carousel.scrollToCenter(e.currentTarget as HTMLElement);
}}
>
<section.icon class="{ICON_CLASS_DEFAULT} flex-shrink-0" />
<span>{section.title}</span>
</button>
{/if}
{/each}
</ScrollCarousel>
</div>
</div>
tools/ui/src/lib/constants/chat-tabs.constants.ts
+18
View File
@@ -0,0 +1,18 @@
/** Sentinel tab id for the bare `#/` new-chat screen */
export const NEW_CHAT_TAB_ID = 'new-chat';
/** Label shown for the new-chat sentinel tab. */
export const NEW_CHAT_LABEL = 'New chat';
/** Fallback label for conversations without an auto-generated title. */
export const UNNAMED_CHAT_LABEL = 'Chat';
/**
* Tab bar max width so it stays clear of the sidebar strip. The expanded strip
* is `md:w-72` and the collapsed one `md:w-12`; these hold the fully tuned
* `max-w-[calc(100vw-?rem)]` classes so the offset has a single source.
*/
export const CHAT_TABS_MAX_WIDTH = {
COLLAPSED_SIDEBAR: 'max-w-[calc(100vw-5rem)]',
EXPANDED_SIDEBAR: 'max-w-[calc(100vw-20rem)]'
} as const;
tools/ui/src/lib/constants/css-classes.constants.ts
+6
View File
@@ -26,5 +26,11 @@ export const CHAT_INPUT_FOCUS_SELECTOR =
/** Default Tailwind size class for inline icon components (lucide, etc.). */
export const ICON_CLASS_DEFAULT = 'h-4 w-4';
/** Small Tailwind size class for inline icons. */
export const ICON_CLASS_SM = 'h-3.5 w-3.5';
/** Extra-small Tailwind size class for inline icons. */
export const ICON_CLASS_XS = 'h-3 w-3';
/** Icon size + spinning animation; used for live-streaming tool indicators. */
export const ICON_CLASS_SPIN = 'h-4 w-4 animate-spin';
tools/ui/src/lib/constants/index.ts
+1
View File
@@ -4,6 +4,7 @@
export * from './agentic.constants';
export * from './api-endpoints.constants';
export * from './app.constants';
export * from './chat-tabs.constants';
export * from './database.constants';
export * from './reasoning-effort.constants';
export * from './recommended-mcp-servers.constants';

Some files were not shown because too many files have changed in this diff Show More