vrr/koboldcpp
2
0
Fork 0
mirror of https://github.com/LostRuins/koboldcpp.git synced 2026-08-10 17:06:13 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

Merge branch 'upstream' into concedo_experimental

Browse source 
# Conflicts:
#	.devops/intel.Dockerfile
#	.github/workflows/build-cross.yml
#	.github/workflows/build-sycl.yml
#	.github/workflows/build.yml
#	.github/workflows/editorconfig.yml
#	.github/workflows/release.yml
#	cmake/riscv64-spacemit-linux-gnu-gcc.cmake
#	docs/backend/OPENVINO.md
#	docs/backend/SYCL.md
#	docs/build-riscv64-spacemit.md
#	docs/ops.md
#	docs/ops/WebGPU.csv
#	embd_res/ggml-vocab-qwen35.gguf
#	embd_res/ggml-vocab-qwen35.gguf.inp
#	embd_res/ggml-vocab-qwen35.gguf.out
#	examples/model-conversion/Makefile
#	ggml/CMakeLists.txt
#	ggml/src/ggml-cpu/CMakeLists.txt
#	ggml/src/ggml-hexagon/ggml-hexagon.cpp
#	ggml/src/ggml-hexagon/htp/hmx-flash-attn-ops.c
#	ggml/src/ggml-hexagon/htp/hmx-matmul-ops.c
#	ggml/src/ggml-hexagon/htp/hmx-utils.h
#	ggml/src/ggml-hexagon/htp/htp-ops.h
#	ggml/src/ggml-hexagon/htp/hvx-utils.h
#	ggml/src/ggml-hexagon/htp/main.c
#	ggml/src/ggml-hexagon/htp/unary-ops.c
#	ggml/src/ggml-opencl/CMakeLists.txt
#	ggml/src/ggml-opencl/ggml-opencl.cpp
#	ggml/src/ggml-opencl/kernels/cvt.cl
#	ggml/src/ggml-sycl/CMakeLists.txt
#	ggml/src/ggml-sycl/common.cpp
#	ggml/src/ggml-sycl/common.hpp
#	ggml/src/ggml-sycl/ggml-sycl.cpp
#	ggml/src/ggml-webgpu/ggml-webgpu-shader-lib.hpp
#	ggml/src/ggml-webgpu/ggml-webgpu.cpp
#	ggml/src/ggml-webgpu/wgsl-shaders/common_decls.tmpl
#	ggml/src/ggml-webgpu/wgsl-shaders/flash_attn_tile.wgsl
#	ggml/src/ggml-webgpu/wgsl-shaders/flash_attn_vec_reduce.wgsl
#	ggml/src/ggml-webgpu/wgsl-shaders/flash_attn_vec_split.wgsl
#	ggml/src/ggml-webgpu/wgsl-shaders/get_rows.wgsl
#	ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_decls.tmpl
#	ggml/src/ggml-webgpu/wgsl-shaders/mul_mat_vec_acc.tmpl
#	ggml/src/ggml-webgpu/wgsl-shaders/unary.wgsl
#	ggml/src/ggml-zendnn/CMakeLists.txt
#	ggml/src/ggml-zendnn/ggml-zendnn.cpp
#	scripts/snapdragon/adb/run-completion.sh
#	tests/CMakeLists.txt
#	tools/cli/README.md
#	tools/completion/README.md
#	tools/mtmd/clip-impl.h
#	tools/mtmd/clip.cpp
#	tools/mtmd/clip.h
#	tools/server/README.md
This commit is contained in:
Concedo 2026-05-14 19:04:04 +08:00
commit cc82c3164e
62 changed files with 18693 additions and 7034 deletions
Download patch file Download diff file Expand all files Collapse all files

66
common/arg.cpp
View file

@ -310,12 +310,14 @@ static bool common_params_handle_remote_preset(common_params & params, llama_exa
common_download_opts opts;
opts.bearer_token = params.hf_token;
opts.offline = params.offline;
LOG_TRC("%s: looking for remote preset at %s\n", __func__, preset_url.c_str());
const int status = common_download_file_single(preset_url, preset_path, opts);
const bool has_preset = status >= 200 && status < 400;
// remote preset is optional, so we don't error out if not found
if (has_preset) {
LOG_INF("applying remote preset from %s\n", preset_url.c_str());
LOG_TRC("%s: applying remote preset from %s\n", __func__, preset_url.c_str());
common_preset_context ctx(ex, /* only_remote_allowed */ true);
common_preset global;
auto remote_presets = ctx.load_from_ini(preset_path, global);
@ -328,7 +330,7 @@ static bool common_params_handle_remote_preset(common_params & params, llama_exa
throw std::runtime_error("Remote preset.ini does not contain [" + std::string(hf_tag) + "] section");
}
} else {
LOG_INF("%s", "no remote preset found, skipping\n");
LOG_TRC("%s: no remote preset found, skipping\n", __func__);
}
return has_preset;
@ -359,8 +361,7 @@ static handle_model_result common_params_handle_model(struct common_params_model
auto download_result = common_download_model(model, opts, true);
if (download_result.model_path.empty()) {
LOG_ERR("error: failed to download model from Hugging Face\n");
exit(1);
throw std::runtime_error("failed to download model from Hugging Face");
}
model.name = model.hf_repo;
@ -382,8 +383,7 @@ static handle_model_result common_params_handle_model(struct common_params_model
opts.offline = offline;
auto download_result = common_download_model(model, opts);
if (download_result.model_path.empty()) {
LOG_ERR("error: failed to download model from %s\n", model.url.c_str());
exit(1);
throw std::runtime_error("failed to download model from " + model.url);
}
}
@ -437,6 +437,25 @@ static bool parse_bool_value(const std::string & value) {
// CLI argument parsing functions
//
void common_params_handle_models(common_params & params, llama_example curr_ex) {
auto res = common_params_handle_model(params.model, params.hf_token, params.offline);
if (params.no_mmproj) {
params.mmproj = {};
} else if (res.found_mmproj && params.mmproj.path.empty() && params.mmproj.url.empty()) {
// optionally, handle mmproj model when -hf is specified
params.mmproj = res.mmproj;
}
// only download mmproj if the current example is using it
for (const auto & ex : mmproj_examples) {
if (curr_ex == ex) {
common_params_handle_model(params.mmproj, params.hf_token, params.offline);
break;
}
}
common_params_handle_model(params.speculative.draft.mparams, params.hf_token, params.offline);
common_params_handle_model(params.vocoder.model, params.hf_token, params.offline);
}
static bool common_params_parse_ex(int argc, char ** argv, common_params_context & ctx_arg) {
common_params & params = ctx_arg.params;
@ -590,22 +609,7 @@ static bool common_params_parse_ex(int argc, char ** argv, common_params_context
// handle model and download
if (!skip_model_download) {
auto res = common_params_handle_model(params.model, params.hf_token, params.offline);
if (params.no_mmproj) {
params.mmproj = {};
} else if (res.found_mmproj && params.mmproj.path.empty() && params.mmproj.url.empty()) {
// optionally, handle mmproj model when -hf is specified
params.mmproj = res.mmproj;
}
// only download mmproj if the current example is using it
for (const auto & ex : mmproj_examples) {
if (ctx_arg.ex == ex) {
common_params_handle_model(params.mmproj, params.hf_token, params.offline);
break;
}
}
common_params_handle_model(params.speculative.draft.mparams, params.hf_token, params.offline);
common_params_handle_model(params.vocoder.model, params.hf_token, params.offline);
common_params_handle_models(params, ctx_arg.ex);
}
// model is required (except for server)
@ -2221,7 +2225,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
if (llama_supports_rpc()) {
add_opt(common_arg(
{"--rpc"}, "SERVERS",
"comma separated list of RPC servers (host:port)",
"comma-separated list of RPC servers (host:port)",
[](common_params & params, const std::string & value) {
add_rpc_devices(value);
GGML_UNUSED(params);
@ -3301,18 +3305,20 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
).set_env("LLAMA_LOG_VERBOSITY"));
add_opt(common_arg(
{"--log-prefix"},
{"--no-log-prefix"},
"Enable prefix in log messages",
[](common_params &) {
common_log_set_prefix(common_log_main(), true);
[](common_params &, bool value) {
common_log_set_prefix(common_log_main(), value);
}
).set_env("LLAMA_LOG_PREFIX"));
).set_env("LLAMA_ARG_LOG_PREFIX"));
add_opt(common_arg(
{"--log-timestamps"},
{"--no-log-timestamps"},
"Enable timestamps in log messages",
[](common_params &) {
common_log_set_timestamps(common_log_main(), true);
[](common_params &, bool value) {
common_log_set_timestamps(common_log_main(), value);
}
).set_env("LLAMA_LOG_TIMESTAMPS"));
).set_env("LLAMA_ARG_LOG_TIMESTAMPS"));
//
// speculative parameters
@ -3553,7 +3559,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPEC_DRAFT_MODEL"));
add_opt(common_arg(
{"--spec-type"}, common_speculative_all_types_str(),
string_format("type of speculative decoding to use when no draft model is provided (default: %s)\n",
string_format("comma-separated list of types of speculative decoding to use (default: %s)\n",
common_speculative_type_name_str(params.speculative.types).c_str()),
[](common_params & params, const std::string & value) {
const auto enabled_types = string_split<std::string>(value, ',');

3
common/arg.h
View file

@ -129,5 +129,8 @@ bool common_params_to_map(int argc, char ** argv, llama_example ex, std::map<com
// see: https://github.com/ggml-org/llama.cpp/issues/18163
void common_params_add_preset_options(std::vector<common_arg> & args);
// Populate model paths (main model, mmproj, etc) from -hf if necessary
void common_params_handle_models(common_params & params, llama_example curr_ex);
// initialize argument parser context - used by test-arg-parser and preset
common_params_context common_params_parser_init(common_params & params, llama_example ex, void(*print_usage)(int, char **) = nullptr);

29
common/common.cpp
View file

@ -372,15 +372,29 @@ void common_init() {
SetConsoleCP(CP_UTF8);
#endif
llama_log_set(common_log_default_callback, NULL);
common_log_set_prefix(common_log_main(), true);
common_log_set_timestamps(common_log_main(), true);
llama_log_set(common_log_default_callback, NULL);
}
void common_params_print_info(const common_params & params) {
#ifdef NDEBUG
const char * build_type = "";
#else
const char * build_type = " (debug)";
#endif
LOG_TRC("%s: build %d (%s) with %s for %s%s\n", __func__, llama_build_number(), llama_commit(), llama_compiler(), llama_build_target(), build_type);
LOG_DBG("build: %d (%s) with %s for %s%s\n", llama_build_number(), llama_commit(), llama_compiler(), llama_build_target(), build_type);
LOG_INF("log_info: verbosity = %d (adjust with the `-lv N` CLI arg)\n", common_log_get_verbosity_thold());
LOG_INF("device_info:\n");
for (size_t i = 0; i < ggml_backend_dev_count(); ++i) {
auto * dev = ggml_backend_dev_get(i);
size_t free, total;
ggml_backend_dev_memory(dev, &free, &total);
LOG_INF(" - %-8s: %s (%zu MiB, %zu MiB free)\n", ggml_backend_dev_name(dev), ggml_backend_dev_description(dev), total / 1024 / 1024, free / 1024 / 1024);
}
LOG_INF("%s\n", common_params_get_system_info(params).c_str());
}
std::string common_params_get_system_info(const common_params & params) {
@ -1153,7 +1167,8 @@ common_init_result::common_init_result(common_params & params) :
auto cparams = common_context_params_to_llama(params);
if (params.fit_params) {
LOG_INF("%s: fitting params to device memory, for bugs during this step try to reproduce them with -fit off, or provide --verbose logs if the bug only occurs with -fit on\n", __func__);
LOG_INF("%s: fitting params to device memory ...\n", __func__);
LOG_INF("%s: (for bugs during this step try to reproduce them with -fit off, or provide --verbose logs if the bug only occurs with -fit on)\n", __func__);
common_fit_params(params.model.path.c_str(), &mparams, &cparams,
params.tensor_split,
params.tensor_buft_overrides.data(),
@ -1202,7 +1217,7 @@ common_init_result::common_init_result(common_params & params) :
// initialize once
for (llama_token i = 0; i < llama_vocab_n_tokens(vocab); i++) {
if (llama_vocab_is_eog(vocab, i)) {
LOG_INF("%s: added %s logit bias = %f\n", __func__, common_token_to_piece(vocab, i).c_str(), -INFINITY);
LOG_TRC("%s: added %s logit bias = %f\n", __func__, common_token_to_piece(vocab, i).c_str(), -INFINITY);
params.sampling.logit_bias_eog.push_back({i, -INFINITY});
}
}
@ -1215,12 +1230,12 @@ common_init_result::common_init_result(common_params & params) :
}
//if (params.sampling.penalty_last_n == -1) {
// LOG_INF("%s: setting penalty_last_n to ctx_size = %d\n", __func__, llama_n_ctx(lctx));
// LOG_TRC("%s: setting penalty_last_n to ctx_size = %d\n", __func__, llama_n_ctx(lctx));
// params.sampling.penalty_last_n = llama_n_ctx(lctx);
//}
//if (params.sampling.dry_penalty_last_n == -1) {
// LOG_INF("%s: setting dry_penalty_last_n to ctx_size = %d\n", __func__, llama_n_ctx(lctx));
// LOG_TRC("%s: setting dry_penalty_last_n to ctx_size = %d\n", __func__, llama_n_ctx(lctx));
// params.sampling.dry_penalty_last_n = llama_n_ctx(lctx);
//}
@ -1428,7 +1443,7 @@ common_context_seq_rm_type common_context_can_seq_rm(llama_context * ctx) {
// try to remove the last tokens
if (!llama_memory_seq_rm(mem, 0, 1, -1)) {
LOG_WRN("%s: the context does not support partial sequence removal\n", __func__);
LOG_TRC("%s: the context does not support partial sequence removal\n", __func__);
res = COMMON_CONTEXT_SEQ_RM_TYPE_FULL;
goto done;
}

8
common/common.h
View file

@ -158,9 +158,9 @@ enum common_params_sampling_config : uint64_t {
enum common_speculative_type {
COMMON_SPECULATIVE_TYPE_NONE, // no speculative decoding
COMMON_SPECULATIVE_TYPE_DRAFT, // draft model
COMMON_SPECULATIVE_TYPE_EAGLE3, // eagle draft model
COMMON_SPECULATIVE_TYPE_NGRAM_SIMPLE, // simple self-speculative decoding
COMMON_SPECULATIVE_TYPE_DRAFT_SIMPLE, // standalone draft model speculative decoding
COMMON_SPECULATIVE_TYPE_DRAFT_EAGLE3, // Eagle3 speculative decoding
COMMON_SPECULATIVE_TYPE_NGRAM_SIMPLE, // simple self-speculative decoding based on n-grams
COMMON_SPECULATIVE_TYPE_NGRAM_MAP_K, // self-speculative decoding with n-gram keys only
COMMON_SPECULATIVE_TYPE_NGRAM_MAP_K4V, // self-speculative decoding with n-gram keys and 4 m-gram values
COMMON_SPECULATIVE_TYPE_NGRAM_MOD,
@ -343,6 +343,7 @@ struct common_params_speculative_ngram_cache {
struct common_params_speculative {
std::vector<enum common_speculative_type> types = { COMMON_SPECULATIVE_TYPE_NONE };
// used by Simple, MTP, Eagle3, etc. - all methods that require some kind of draft model
common_params_speculative_draft draft;
common_params_speculative_ngram_mod ngram_mod;
@ -686,6 +687,7 @@ struct common_params {
// initializes the logging system and prints info about the build
void common_init();
void common_params_print_info(const common_params & params);
std::string common_params_get_system_info(const common_params & params);
bool parse_cpu_range(const std::string & range, bool(&boolmask)[GGML_MAX_N_THREADS]);

4
common/download.cpp
View file

@ -323,9 +323,9 @@ static int common_download_file_single_online(const std::string & url,
auto head = cli.Head(parts.path);
if (!head || head->status < 200 || head->status >= 300) {
LOG_WRN("%s: HEAD failed, status: %d\n", __func__, head ? head->status : -1);
LOG_TRC("%s: HEAD failed, status: %d\n", __func__, head ? head->status : -1);
if (file_exists) {
LOG_INF("%s: using cached file (HEAD failed): %s\n", __func__, path.c_str());
LOG_TRC("%s: using cached file (HEAD failed): %s\n", __func__, path.c_str());
return 304; // 304 Not Modified - fake cached response
}
return head ? head->status : -1;

88
common/fit.cpp
View file

@ -168,7 +168,7 @@ static void common_params_fit_impl(
// step 1: get data for default parameters and check whether any changes are necessary in the first place
LOG_INF("%s: getting device memory data for initial parameters:\n", __func__);
LOG_TRC("%s: getting device memory data for initial parameters:\n", __func__);
const dmds_t dmds_full = common_get_device_memory_data(path_model, mparams, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level);
const size_t nd = devs.size(); // number of devices
@ -213,13 +213,13 @@ static void common_params_fit_impl(
LOG_INF("%s: projected to use %" PRId64 " MiB of host memory vs. %" PRId64 " MiB of total host memory\n",
__func__, sum_projected_used/MiB, sum_free/MiB);
if (sum_projected_free >= margins[0]) {
LOG_INF("%s: will leave %" PRId64 " >= %" PRId64 " MiB of system memory, no changes needed\n",
LOG_TRC("%s: will leave %" PRId64 " >= %" PRId64 " MiB of system memory, no changes needed\n",
__func__, sum_projected_free/MiB, margins[0]/MiB);
return;
}
} else {
if (nd > 1) {
LOG_INF("%s: projected memory use with initial parameters [MiB]:\n", __func__);
LOG_TRC("%s: projected memory use with initial parameters [MiB]:\n", __func__);
}
for (size_t id = 0; id < nd; id++) {
const llama_device_memory_data & dmd = dmds_full[id];
@ -234,16 +234,16 @@ static void common_params_fit_impl(
sum_projected_model += dmd.mb.model;
if (nd > 1) {
LOG_INF("%s: - %s: %6" PRId64 " total, %6" PRId64 " used, %6" PRId64 " free vs. target of %6" PRId64 "\n",
LOG_TRC("%s: - %s: %6" PRId64 " total, %6" PRId64 " used, %6" PRId64 " free vs. target of %6" PRId64 "\n",
__func__, dev_names[id].c_str(), dmd.total/MiB, projected_used/MiB, projected_free/MiB, margins[id]/MiB);
}
}
assert(sum_free >= 0 && sum_projected_used >= 0);
LOG_INF("%s: projected to use %" PRId64 " MiB of device memory vs. %" PRId64 " MiB of free device memory\n",
LOG_TRC("%s: projected to use %" PRId64 " MiB of device memory vs. %" PRId64 " MiB of free device memory\n",
__func__, sum_projected_used/MiB, sum_free/MiB);
if (nd == 1) {
if (projected_free_per_device[0] >= margins[0]) {
LOG_INF("%s: will leave %" PRId64 " >= %" PRId64 " MiB of free device memory, no changes needed\n",
LOG_TRC("%s: will leave %" PRId64 " >= %" PRId64 " MiB of free device memory, no changes needed\n",
__func__, projected_free_per_device[0]/MiB, margins[0]/MiB);
return;
}
@ -256,7 +256,7 @@ static void common_params_fit_impl(
}
}
if (!changes_needed) {
LOG_INF("%s: targets for free memory can be met on all devices, no changes needed\n", __func__);
LOG_TRC("%s: targets for free memory can be met on all devices, no changes needed\n", __func__);
return;
}
}
@ -275,10 +275,10 @@ static void common_params_fit_impl(
}
if (global_surplus < 0) {
if (nd <= 1) {
LOG_INF("%s: cannot meet free memory target of %" PRId64 " MiB, need to reduce device memory by %" PRId64 " MiB\n",
LOG_TRC("%s: cannot meet free memory target of %" PRId64 " MiB, need to reduce device memory by %" PRId64 " MiB\n",
__func__, margins[0]/MiB, -global_surplus/MiB);
} else {
LOG_INF(
LOG_TRC(
"%s: cannot meet free memory targets on all devices, need to use %" PRId64 " MiB less in total\n",
__func__, -global_surplus/MiB);
}
@ -320,28 +320,28 @@ static void common_params_fit_impl(
const int64_t bytes_per_ctx = (sum_projected_used - sum_projected_used_min_ctx) / (hp_nct - n_ctx_min);
const int64_t memory_reduction = (hp_nct - cparams->n_ctx) * bytes_per_ctx;
LOG_INF("%s: context size reduced from %" PRIu32 " to %" PRIu32 " -> need %" PRId64 " MiB less memory in total\n",
LOG_TRC("%s: context size reduced from %" PRIu32 " to %" PRIu32 " -> need %" PRId64 " MiB less memory in total\n",
__func__, hp_nct, cparams->n_ctx, memory_reduction/MiB);
if (nd <= 1) {
LOG_INF("%s: entire model can be fit by reducing context\n", __func__);
LOG_TRC("%s: entire model can be fit by reducing context\n", __func__);
return;
}
LOG_INF("%s: entire model should be fit across devices by reducing context\n", __func__);
LOG_TRC("%s: entire model should be fit across devices by reducing context\n", __func__);
} else {
const int64_t memory_reduction = sum_projected_used - sum_projected_used_min_ctx;
LOG_INF("%s: context size reduced from %" PRIu32 " to %" PRIu32 " -> need %" PRId64 " MiB less memory in total\n",
LOG_TRC("%s: context size reduced from %" PRIu32 " to %" PRIu32 " -> need %" PRId64 " MiB less memory in total\n",
__func__, hp_nct, cparams->n_ctx, memory_reduction/MiB);
}
} else {
if (n_ctx_min == UINT32_MAX) {
LOG_INF("%s: user has requested full context size of %" PRIu32 " -> no change\n", __func__, hp_nct);
LOG_TRC("%s: user has requested full context size of %" PRIu32 " -> no change\n", __func__, hp_nct);
} else {
LOG_INF("%s: default model context size is %" PRIu32 " which is <= the min. context size of %" PRIu32 " -> no change\n",
LOG_TRC("%s: default model context size is %" PRIu32 " which is <= the min. context size of %" PRIu32 " -> no change\n",
__func__, hp_nct, n_ctx_min);
}
}
} else {
LOG_INF("%s: context size set by user to %" PRIu32 " -> no change\n", __func__, cparams->n_ctx);
LOG_TRC("%s: context size set by user to %" PRIu32 " -> no change\n", __func__, cparams->n_ctx);
}
}
}
@ -485,10 +485,10 @@ static void common_params_fit_impl(
const dmds_t dmd_nl = common_get_device_memory_data(
path_model, &mparams_copy, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level);
LOG_INF("%s: memory for test allocation by device:\n", func_name);
LOG_TRC("%s: memory for test allocation by device:\n", func_name);
for (size_t id = 0; id < nd; id++) {
const ngl_t & n = ngl_per_device[id];
LOG_INF(
LOG_TRC(
"%s: id=%zu, n_layer=%2" PRIu32 ", n_part=%2" PRIu32 ", overflow_type=%d, mem=%6" PRId64 " MiB\n",
func_name, id, n.n_layer, n.n_part, int(n.overflow_type), dmd_nl[id].mb.total()/MiB);
}
@ -509,7 +509,7 @@ static void common_params_fit_impl(
tensor_buft_overrides[1] = {nullptr, nullptr};
mparams->tensor_buft_overrides = tensor_buft_overrides;
LOG_INF("%s: getting device memory data with all MoE tensors moved to system memory:\n", __func__);
LOG_TRC("%s: getting device memory data with all MoE tensors moved to system memory:\n", __func__);
const dmds_t dmds_cpu_moe = common_get_device_memory_data(
path_model, mparams, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level);
@ -519,10 +519,10 @@ static void common_params_fit_impl(
}
if (global_surplus_cpu_moe > 0) {
LOG_INF("%s: with only dense weights in device memory there is a total surplus of %" PRId64 " MiB\n",
LOG_TRC("%s: with only dense weights in device memory there is a total surplus of %" PRId64 " MiB\n",
__func__, global_surplus_cpu_moe/MiB);
} else {
LOG_INF("%s: with only dense weights in device memory there is still a total deficit of %" PRId64 " MiB\n",
LOG_TRC("%s: with only dense weights in device memory there is still a total deficit of %" PRId64 " MiB\n",
__func__, -global_surplus_cpu_moe/MiB);
}
@ -535,7 +535,7 @@ static void common_params_fit_impl(
targets.reserve(nd);
for (size_t id = 0; id < nd; id++) {
targets.push_back(dmds_full[id].free - margins[id]);
LOG_INF("%s: id=%zu, target=%" PRId64 " MiB\n", __func__, id, targets[id]/MiB);
LOG_TRC("%s: id=%zu, target=%" PRId64 " MiB\n", __func__, id, targets[id]/MiB);
}
std::vector<ggml_backend_buffer_type_t> overflow_bufts; // which bufts the first partial layer of a device overflows to:
@ -555,9 +555,9 @@ static void common_params_fit_impl(
// - once we only have a difference of a single layer, stop and return the lower bound that just barely still fits
// - the last device has the output layer, which cannot be a partial layer
if (hp_nex == 0) {
LOG_INF("%s: filling dense layers back-to-front:\n", __func__);
LOG_TRC("%s: filling dense layers back-to-front:\n", __func__);
} else {
LOG_INF("%s: filling dense-only layers back-to-front:\n", __func__);
LOG_TRC("%s: filling dense-only layers back-to-front:\n", __func__);
}
for (int id = nd - 1; id >= 0; id--) {
uint32_t n_unassigned = hp_ngl + 1;
@ -576,7 +576,7 @@ static void common_params_fit_impl(
if (mem_high[id] > targets[id]) {
assert(ngl_per_device_high[id].n_layer > ngl_per_device[id].n_layer);
uint32_t delta = ngl_per_device_high[id].n_layer - ngl_per_device[id].n_layer;
LOG_INF("%s: start filling device %" PRIu32 ", delta=%" PRIu32 "\n", __func__, id, delta);
LOG_TRC("%s: start filling device %" PRIu32 ", delta=%" PRIu32 "\n", __func__, id, delta);
while (delta > 1) {
uint32_t step_size = int64_t(delta) * (targets[id] - mem[id]) / (mem_high[id] - mem[id]);
step_size = std::max(step_size, uint32_t(1));
@ -593,11 +593,11 @@ static void common_params_fit_impl(
if (mem_test[id] <= targets[id]) {
ngl_per_device = ngl_per_device_test;
mem = mem_test;
LOG_INF("%s: set ngl_per_device[%d].n_layer=%" PRIu32 "\n", __func__, id, ngl_per_device[id].n_layer);
LOG_TRC("%s: set ngl_per_device[%d].n_layer=%" PRIu32 "\n", __func__, id, ngl_per_device[id].n_layer);
} else {
ngl_per_device_high = ngl_per_device_test;
mem_high = mem_test;
LOG_INF("%s: set ngl_per_device_high[%d].n_layer=%" PRIu32 "\n", __func__, id, ngl_per_device_high[id].n_layer);
LOG_TRC("%s: set ngl_per_device_high[%d].n_layer=%" PRIu32 "\n", __func__, id, ngl_per_device_high[id].n_layer);
}
delta = ngl_per_device_high[id].n_layer - ngl_per_device[id].n_layer;
}
@ -605,12 +605,12 @@ static void common_params_fit_impl(
assert(ngl_per_device_high[id].n_layer == n_unassigned);
ngl_per_device = ngl_per_device_high;
mem = mem_high;
LOG_INF("%s: set ngl_per_device[%d].n_layer=%" PRIu32 "\n", __func__, id, ngl_per_device[id].n_layer);
LOG_TRC("%s: set ngl_per_device[%d].n_layer=%" PRIu32 "\n", __func__, id, ngl_per_device[id].n_layer);
}
}
const int64_t projected_margin = dmds_full[id].free - mem[id];
LOG_INF(
LOG_TRC(
"%s: - %s: %2" PRIu32 " layers, %6" PRId64 " MiB used, %6" PRId64 " MiB free\n",
__func__, dev_names[id].c_str(), ngl_per_device[id].n_layer, mem[id]/MiB, projected_margin/MiB);
}
@ -634,7 +634,7 @@ static void common_params_fit_impl(
}
assert(id_dense_start < nd);
LOG_INF("%s: converting dense-only layers to full layers and filling them front-to-back with overflow to next device/system memory:\n", __func__);
LOG_TRC("%s: converting dense-only layers to full layers and filling them front-to-back with overflow to next device/system memory:\n", __func__);
for (size_t id = 0; id <= id_dense_start && id_dense_start < nd; id++) {
std::vector<ngl_t> ngl_per_device_high = ngl_per_device;
for (size_t jd = id_dense_start; jd < nd; jd++) {
@ -674,13 +674,13 @@ static void common_params_fit_impl(
ngl_per_device = ngl_per_device_test;
mem = mem_test;
id_dense_start = id_dense_start_test;
LOG_INF("%s: set ngl_per_device[%zu].(n_layer, n_part)=(%" PRIu32 ", %" PRIu32 "), id_dense_start=%zu\n",
LOG_TRC("%s: set ngl_per_device[%zu].(n_layer, n_part)=(%" PRIu32 ", %" PRIu32 "), id_dense_start=%zu\n",
__func__, id, ngl_per_device[id].n_layer, ngl_per_device[id].n_part, id_dense_start);
} else {
ngl_per_device_high = ngl_per_device_test;
mem_high = mem_test;
id_dense_start_high = id_dense_start_test;
LOG_INF("%s: set ngl_per_device_high[%zu].(n_layer, n_part)=(%" PRIu32 ", %" PRIu32 "), id_dense_start_high=%zu\n",
LOG_TRC("%s: set ngl_per_device_high[%zu].(n_layer, n_part)=(%" PRIu32 ", %" PRIu32 "), id_dense_start_high=%zu\n",
__func__, id, ngl_per_device_high[id].n_layer, ngl_per_device_high[id].n_part, id_dense_start_high);
}
assert(ngl_per_device_high[id].n_full() >= ngl_per_device[id].n_full());
@ -690,7 +690,7 @@ static void common_params_fit_impl(
ngl_per_device = ngl_per_device_high;
mem = mem_high;
id_dense_start = id_dense_start_high;
LOG_INF("%s: set ngl_per_device[%zu].(n_layer, n_part)=(%" PRIu32 ", %" PRIu32 "), id_dense_start=%zu\n",
LOG_TRC("%s: set ngl_per_device[%zu].(n_layer, n_part)=(%" PRIu32 ", %" PRIu32 "), id_dense_start=%zu\n",
__func__, id, ngl_per_device[id].n_layer, ngl_per_device[id].n_part, id_dense_start);
}
@ -710,44 +710,44 @@ static void common_params_fit_impl(
if (id < nd - 1) {
overflow_bufts_test[id] = ggml_backend_dev_buffer_type(devs[id + 1]);
}
LOG_INF("%s: trying to fit one extra layer with overflow_type=LAYER_FRACTION_UP\n", __func__);
LOG_TRC("%s: trying to fit one extra layer with overflow_type=LAYER_FRACTION_UP\n", __func__);
std::vector<int64_t> mem_test = get_memory_for_layers(__func__, ngl_per_device_test, overflow_bufts_test);
if (mem_test[id] < targets[id] && (id + 1 == nd || mem_test[id + 1] < targets[id + 1])) {
ngl_per_device = ngl_per_device_test;
overflow_bufts = overflow_bufts_test;
mem = mem_test;
id_dense_start = id_dense_start_test;
LOG_INF("%s: set ngl_per_device[%zu].(n_layer, n_part, overflow_type)=(%" PRIu32 ", %" PRIu32 ", UP), id_dense_start=%zu\n",
LOG_TRC("%s: set ngl_per_device[%zu].(n_layer, n_part, overflow_type)=(%" PRIu32 ", %" PRIu32 ", UP), id_dense_start=%zu\n",
__func__, id, ngl_per_device[id].n_layer, ngl_per_device[id].n_part, id_dense_start);
ngl_per_device_test[id].overflow_type = LAYER_FRACTION_GATE;
LOG_INF("%s: trying to fit one extra layer with overflow_type=LAYER_FRACTION_GATE\n", __func__);
LOG_TRC("%s: trying to fit one extra layer with overflow_type=LAYER_FRACTION_GATE\n", __func__);
mem_test = get_memory_for_layers(__func__, ngl_per_device_test, overflow_bufts_test);
if (mem_test[id] < targets[id] && (id + 1 == nd || mem_test[id + 1] < targets[id + 1])) {
ngl_per_device = ngl_per_device_test;
overflow_bufts = overflow_bufts_test;
mem = mem_test;
id_dense_start = id_dense_start_test;
LOG_INF("%s: set ngl_per_device[%zu].(n_layer, n_part, overflow_type)=(%" PRIu32 ", %" PRIu32 ", GATE), id_dense_start=%zu\n",
LOG_TRC("%s: set ngl_per_device[%zu].(n_layer, n_part, overflow_type)=(%" PRIu32 ", %" PRIu32 ", GATE), id_dense_start=%zu\n",
__func__, id, ngl_per_device[id].n_layer, ngl_per_device[id].n_part, id_dense_start);
}
} else {
ngl_per_device_test[id].overflow_type = LAYER_FRACTION_ATTN;
LOG_INF("%s: trying to fit one extra layer with overflow_type=LAYER_FRACTION_ATTN\n", __func__);
LOG_TRC("%s: trying to fit one extra layer with overflow_type=LAYER_FRACTION_ATTN\n", __func__);
mem_test = get_memory_for_layers(__func__, ngl_per_device_test, overflow_bufts_test);
if (mem_test[id] < targets[id] && (id + 1 == nd || mem_test[id + 1] < targets[id + 1])) {
ngl_per_device = ngl_per_device_test;
overflow_bufts = overflow_bufts_test;
mem = mem_test;
id_dense_start = id_dense_start_test;
LOG_INF("%s: set ngl_per_device[%zu].(n_layer, n_part, overflow_type)=(%" PRIu32 ", %" PRIu32 ", ATTN), id_dense_start=%zu\n",
LOG_TRC("%s: set ngl_per_device[%zu].(n_layer, n_part, overflow_type)=(%" PRIu32 ", %" PRIu32 ", ATTN), id_dense_start=%zu\n",
__func__, id, ngl_per_device[id].n_layer, ngl_per_device[id].n_part, id_dense_start);
}
}
}
const int64_t projected_margin = dmds_full[id].free - mem[id];
LOG_INF(
LOG_TRC(
"%s: - %s: %2" PRIu32 " layers (%2" PRIu32 " overflowing), %6" PRId64 " MiB used, %6" PRId64 " MiB free\n",
__func__, dev_names[id].c_str(), ngl_per_device[id].n_layer, ngl_per_device[id].n_part, mem[id]/MiB, projected_margin/MiB);
}
@ -755,7 +755,7 @@ static void common_params_fit_impl(
// print info for devices that were not changed during the conversion from dense only to full layers:
for (size_t id = id_dense_start + 1; id < nd; id++) {
const int64_t projected_margin = dmds_full[id].free - mem[id];
LOG_INF(
LOG_TRC(
"%s: - %s: %2" PRIu32 " layers (%2" PRIu32 " overflowing), %6" PRId64 " MiB used, %6" PRId64 " MiB free\n",
__func__, dev_names[id].c_str(), ngl_per_device[id].n_layer, ngl_per_device[id].n_part, mem[id]/MiB, projected_margin/MiB);
}
@ -776,7 +776,7 @@ enum common_params_fit_status common_fit_params(
common_params_fit_status status = COMMON_PARAMS_FIT_STATUS_SUCCESS;
try {
common_params_fit_impl(path_model, mparams, cparams, tensor_split, tensor_buft_overrides, margins, n_ctx_min, log_level);
LOG_INF("%s: successfully fit params to free device memory\n", __func__);
LOG_TRC("%s: successfully fit params to free device memory\n", __func__);
} catch (const common_params_fit_exception & e) {
LOG_WRN("%s: failed to fit params to free device memory: %s\n", __func__, e.what());
status = COMMON_PARAMS_FIT_STATUS_FAILURE;
@ -785,7 +785,7 @@ enum common_params_fit_status common_fit_params(
status = COMMON_PARAMS_FIT_STATUS_ERROR;
}
const int64_t t1_us = llama_time_us();
LOG_INF("%s: fitting params to free memory took %.2f seconds\n", __func__, (t1_us - t0_us) * 1e-6);
LOG_TRC("%s: fitting params to free memory took %.2f seconds\n", __func__, (t1_us - t0_us) * 1e-6);
return status;
}
@ -925,7 +925,7 @@ void common_memory_breakdown_print(const struct llama_context * ctx) {
}
}
for (const auto & td : table_data) {
LOG_INF(td[0].c_str(),
LOG_TRC(td[0].c_str(),
__func__, td[1].c_str(), td[2].c_str(), td[3].c_str(), td[4].c_str(), td[5].c_str(),
td[6].c_str(), td[7].c_str(), td[8].c_str());
}

4
common/log.cpp
View file

@ -435,10 +435,10 @@ void common_log_flush(struct common_log * log) {
static int common_get_verbosity(enum ggml_log_level level) {
switch (level) {
case GGML_LOG_LEVEL_DEBUG: return LOG_LEVEL_DEBUG;
case GGML_LOG_LEVEL_INFO: return LOG_LEVEL_INFO;
case GGML_LOG_LEVEL_INFO: return LOG_LEVEL_TRACE;
case GGML_LOG_LEVEL_WARN: return LOG_LEVEL_WARN;
case GGML_LOG_LEVEL_ERROR: return LOG_LEVEL_ERROR;
case GGML_LOG_LEVEL_CONT: return LOG_LEVEL_INFO; // same as INFO
case GGML_LOG_LEVEL_CONT: return LOG_LEVEL_TRACE;
case GGML_LOG_LEVEL_NONE:
default:
return LOG_LEVEL_OUTPUT;

7
common/log.h
View file

@ -21,7 +21,8 @@
# define LOG_ATTRIBUTE_FORMAT(...) __attribute__((format(printf, __VA_ARGS__)))
#endif
#define LOG_LEVEL_DEBUG 4
#define LOG_LEVEL_DEBUG 5
#define LOG_LEVEL_TRACE 4
#define LOG_LEVEL_INFO 3
#define LOG_LEVEL_WARN 2
#define LOG_LEVEL_ERROR 1
@ -111,13 +112,15 @@ void common_log_flush (struct common_log * log); // f
#define LOGV(verbosity, ...) LOG_TMPL(GGML_LOG_LEVEL_NONE, verbosity, __VA_ARGS__)
#define LOG_DBG(...) LOG_TMPL(GGML_LOG_LEVEL_DEBUG, LOG_LEVEL_DEBUG, __VA_ARGS__)
#define LOG_TRC(...) LOG_TMPL(GGML_LOG_LEVEL_INFO, LOG_LEVEL_TRACE, __VA_ARGS__)
#define LOG_INF(...) LOG_TMPL(GGML_LOG_LEVEL_INFO, LOG_LEVEL_INFO, __VA_ARGS__)
#define LOG_WRN(...) LOG_TMPL(GGML_LOG_LEVEL_WARN, LOG_LEVEL_WARN, __VA_ARGS__)
#define LOG_ERR(...) LOG_TMPL(GGML_LOG_LEVEL_ERROR, LOG_LEVEL_ERROR, __VA_ARGS__)
#define LOG_CNT(...) LOG_TMPL(GGML_LOG_LEVEL_CONT, LOG_LEVEL_INFO, __VA_ARGS__) // same as INFO
#define LOG_DBGV(verbosity, ...) LOG_TMPL(GGML_LOG_LEVEL_DEBUG, verbosity, __VA_ARGS__)
#define LOG_TRCV(verbosity, ...) LOG_TMPL(GGML_LOG_LEVEL_TRACE, verbosity, __VA_ARGS__)
#define LOG_INFV(verbosity, ...) LOG_TMPL(GGML_LOG_LEVEL_INFO, verbosity, __VA_ARGS__)
#define LOG_WRNV(verbosity, ...) LOG_TMPL(GGML_LOG_LEVEL_WARN, verbosity, __VA_ARGS__)
#define LOG_ERRV(verbosity, ...) LOG_TMPL(GGML_LOG_LEVEL_ERROR, verbosity, __VA_ARGS__)
#define LOG_DBGV(verbosity, ...) LOG_TMPL(GGML_LOG_LEVEL_DEBUG, verbosity, __VA_ARGS__)
#define LOG_CNTV(verbosity, ...) LOG_TMPL(GGML_LOG_LEVEL_CONT, verbosity, __VA_ARGS__)

7
common/preset.cpp
View file

@ -163,8 +163,13 @@ void common_preset::merge(const common_preset & other) {
}
}
void common_preset::apply_to_params(common_params & params) const {
void common_preset::apply_to_params(common_params & params, const std::set<std::string> & handled_keys) const {
for (const auto & [opt, val] : options) {
if (!handled_keys.empty()) {
if (!opt.env || handled_keys.find(opt.env) == handled_keys.end()) {
continue;
}
}
// apply each option to params
if (opt.handler_string) {
opt.handler_string(params, val);

3
common/preset.h
View file

@ -43,7 +43,8 @@ struct common_preset {
void merge(const common_preset & other);
// apply preset options to common_params
void apply_to_params(common_params & params) const;
// optionally specify handled_keys to only apply a subset of options (identified by their env), if empty, apply all options
void apply_to_params(common_params & params, const std::set<std::string> & handled_keys = std::set<std::string>()) const;
};
// interface for multiple presets in one file

80
common/speculative.cpp
View file

@ -21,8 +21,8 @@
const std::map<std::string, common_speculative_type> common_speculative_type_from_name_map = {
{"none", COMMON_SPECULATIVE_TYPE_NONE},
{"draft", COMMON_SPECULATIVE_TYPE_DRAFT},
{"eagle3", COMMON_SPECULATIVE_TYPE_EAGLE3},
{"draft-simple", COMMON_SPECULATIVE_TYPE_DRAFT_SIMPLE},
{"draft-eagle3", COMMON_SPECULATIVE_TYPE_DRAFT_EAGLE3},
{"ngram-simple", COMMON_SPECULATIVE_TYPE_NGRAM_SIMPLE},
{"ngram-map-k", COMMON_SPECULATIVE_TYPE_NGRAM_MAP_K},
{"ngram-map-k4v", COMMON_SPECULATIVE_TYPE_NGRAM_MAP_K4V},
@ -145,15 +145,15 @@ struct common_speculative_impl {
virtual void accept(llama_seq_id seq_id, uint16_t n_accepted) = 0;
};
struct common_speculative_state_draft : public common_speculative_impl {
struct common_speculative_impl_draft_simple : public common_speculative_impl {
common_params_speculative_draft params;
llama_batch batch;
std::vector<common_sampler_ptr> smpls;
common_speculative_state_draft(const common_params_speculative & params, uint32_t n_seq)
: common_speculative_impl(COMMON_SPECULATIVE_TYPE_DRAFT, n_seq)
common_speculative_impl_draft_simple(const common_params_speculative & params, uint32_t n_seq)
: common_speculative_impl(COMMON_SPECULATIVE_TYPE_DRAFT_SIMPLE, n_seq)
, params(params.draft)
{
auto * ctx_dft = this->params.ctx_dft;
@ -206,7 +206,7 @@ struct common_speculative_state_draft : public common_speculative_impl {
}
}
~common_speculative_state_draft() override {
~common_speculative_impl_draft_simple() override {
llama_batch_free(batch);
}
@ -340,11 +340,11 @@ struct common_speculative_state_draft : public common_speculative_impl {
}
};
struct common_speculative_state_eagle3 : public common_speculative_impl {
struct common_speculative_impl_draft_eagle3 : public common_speculative_impl {
//common_params_speculative_eagle3 params;
common_speculative_state_eagle3(const common_params_speculative & /*params*/, uint32_t n_seq)
: common_speculative_impl(COMMON_SPECULATIVE_TYPE_EAGLE3, n_seq) {}
common_speculative_impl_draft_eagle3(const common_params_speculative & /*params*/, uint32_t n_seq)
: common_speculative_impl(COMMON_SPECULATIVE_TYPE_DRAFT_EAGLE3, n_seq) {}
void begin(llama_seq_id /*seq_id*/, const llama_tokens & /*prompt*/) override {
// noop
@ -365,13 +365,13 @@ struct common_speculative_state_eagle3 : public common_speculative_impl {
};
// state of self-speculation (simple implementation, not ngram-map)
struct common_speculative_state_ngram_simple : public common_speculative_impl {
struct common_speculative_impl_ngram_simple : public common_speculative_impl {
common_params_speculative_ngram_map params;
// shared across all sequences
common_ngram_simple_config config;
common_speculative_state_ngram_simple(
common_speculative_impl_ngram_simple(
const common_params_speculative & params, uint32_t n_seq,
common_ngram_simple_config config)
: common_speculative_impl(COMMON_SPECULATIVE_TYPE_NGRAM_SIMPLE, n_seq)
@ -405,13 +405,13 @@ struct common_speculative_state_ngram_simple : public common_speculative_impl {
}
};
struct common_speculative_state_ngram_map_k : public common_speculative_impl {
struct common_speculative_impl_ngram_map_k : public common_speculative_impl {
common_params_speculative_ngram_map params;
// n_seq configs
std::vector<common_ngram_map> config;
common_speculative_state_ngram_map_k(
common_speculative_impl_ngram_map_k(
const common_params_speculative & params,
const common_ngram_map & config,
uint32_t n_seq)
@ -453,7 +453,7 @@ struct common_speculative_state_ngram_map_k : public common_speculative_impl {
}
};
struct common_speculative_state_ngram_mod : public common_speculative_impl {
struct common_speculative_impl_ngram_mod : public common_speculative_impl {
common_params_speculative_ngram_mod params;
// shared across all sequences
@ -475,7 +475,7 @@ struct common_speculative_state_ngram_mod : public common_speculative_impl {
std::vector<seq_info> sinfos;
common_speculative_state_ngram_mod(
common_speculative_impl_ngram_mod(
const common_params_speculative & params,
uint32_t n_seq)
: common_speculative_impl(COMMON_SPECULATIVE_TYPE_NGRAM_MOD, n_seq)
@ -621,7 +621,7 @@ struct common_speculative_state_ngram_mod : public common_speculative_impl {
}
};
struct common_speculative_state_ngram_cache : public common_speculative_impl {
struct common_speculative_impl_ngram_cache : public common_speculative_impl {
common_params_speculative_ngram_cache params;
uint16_t n_draft;
@ -639,7 +639,7 @@ struct common_speculative_state_ngram_cache : public common_speculative_impl {
std::vector<seq_info> sinfos;
common_speculative_state_ngram_cache(
common_speculative_impl_ngram_cache(
const common_params_speculative & params,
uint32_t n_seq,
uint16_t n_draft,
@ -775,7 +775,7 @@ static common_ngram_map get_common_ngram_map(
return common_ngram_map(size_key, size_value, key_only, min_hits);
}
static common_speculative_state_ngram_cache create_state_ngram_cache(
static common_speculative_impl_ngram_cache create_state_ngram_cache(
const common_speculative_config & config,
uint32_t n_seq,
const std::string & path_static,
@ -786,7 +786,7 @@ static common_speculative_state_ngram_cache create_state_ngram_cache(
bool save_static = false;
bool save_dynamic = false;
common_speculative_state_ngram_cache state(config.params, n_seq, n_draft, path_static, path_dynamic, save_static, save_dynamic);
common_speculative_impl_ngram_cache state(config.params, n_seq, n_draft, path_static, path_dynamic, save_static, save_dynamic);
return state;
}
@ -818,8 +818,8 @@ const char * common_speculative_all_types_str() {
std::string common_speculative_type_to_str(common_speculative_type type) {
switch (type) {
case COMMON_SPECULATIVE_TYPE_NONE: return "none";
case COMMON_SPECULATIVE_TYPE_DRAFT: return "draft";
case COMMON_SPECULATIVE_TYPE_EAGLE3: return "eagle3";
case COMMON_SPECULATIVE_TYPE_DRAFT_SIMPLE: return "draft-simple";
case COMMON_SPECULATIVE_TYPE_DRAFT_EAGLE3: return "draft-eagle3";
case COMMON_SPECULATIVE_TYPE_NGRAM_SIMPLE: return "ngram-simple";
case COMMON_SPECULATIVE_TYPE_NGRAM_MAP_K: return "ngram-map-k";
case COMMON_SPECULATIVE_TYPE_NGRAM_MAP_K4V: return "ngram-map-k4v";
@ -872,9 +872,9 @@ common_speculative * common_speculative_init(common_params_speculative & params,
{
uint32_t enabled_configs = common_get_enabled_speculative_configs(params.types);
bool has_draft = (enabled_configs & (1u << COMMON_SPECULATIVE_TYPE_DRAFT));
bool has_draft_model = !params.draft.mparams.path.empty();
bool has_draft_model_path = !params.draft.mparams.path.empty();
bool has_draft_simple = (enabled_configs & (1u << COMMON_SPECULATIVE_TYPE_DRAFT_SIMPLE));
// bool has_mtp = false; // TODO: add MTP here
bool has_draft_eagle3 = false; // TODO PR-18039: if params.speculative.eagle3
@ -906,22 +906,22 @@ common_speculative * common_speculative_init(common_params_speculative & params,
if (has_ngram_cache) {
configs.push_back(common_speculative_config(COMMON_SPECULATIVE_TYPE_NGRAM_CACHE, params));
}
if (has_draft) {
if (!has_draft_model) {
if (has_draft_simple) {
if (!has_draft_model_path) {
LOG_WRN("%s: draft model is not specified - cannot use 'draft' type\n", __func__);
has_draft = false;
has_draft_simple = false;
}
} else if (has_draft_model) {
} else if (has_draft_model_path) {
LOG_WRN("%s: draft model is specified but 'draft' speculative type is not explicitly enabled - enabling it\n", __func__);
has_draft = true;
has_draft_simple = true;
}
if (has_draft) {
configs.push_back(common_speculative_config(COMMON_SPECULATIVE_TYPE_DRAFT, params));
if (has_draft_simple) {
configs.push_back(common_speculative_config(COMMON_SPECULATIVE_TYPE_DRAFT_SIMPLE, params));
}
// TODO: add MTP here
if (has_draft_eagle3) {
configs.push_back(common_speculative_config(COMMON_SPECULATIVE_TYPE_EAGLE3, params));
configs.push_back(common_speculative_config(COMMON_SPECULATIVE_TYPE_DRAFT_EAGLE3, params));
}
}
@ -932,12 +932,12 @@ common_speculative * common_speculative_init(common_params_speculative & params,
switch (config.type) {
case COMMON_SPECULATIVE_TYPE_NONE:
break;
case COMMON_SPECULATIVE_TYPE_DRAFT: {
impls.push_back(std::make_unique<common_speculative_state_draft>(config.params, n_seq));
case COMMON_SPECULATIVE_TYPE_DRAFT_SIMPLE: {
impls.push_back(std::make_unique<common_speculative_impl_draft_simple>(config.params, n_seq));
break;
}
case COMMON_SPECULATIVE_TYPE_EAGLE3: {
impls.push_back(std::make_unique<common_speculative_state_eagle3>(config.params, n_seq));
case COMMON_SPECULATIVE_TYPE_DRAFT_EAGLE3: {
impls.push_back(std::make_unique<common_speculative_impl_draft_eagle3>(config.params, n_seq));
break;
}
case COMMON_SPECULATIVE_TYPE_NGRAM_SIMPLE: {
@ -950,7 +950,7 @@ common_speculative * common_speculative_init(common_params_speculative & params,
/* .size_ngram = */ ngram_size_key,
/* .size_mgram = */ mgram_size_value
};
auto state = std::make_unique<common_speculative_state_ngram_simple>(
auto state = std::make_unique<common_speculative_impl_ngram_simple>(
/* .params = */ config.params,
/* .n_seq = */ n_seq,
/* .state = */ config_simple
@ -961,13 +961,13 @@ common_speculative * common_speculative_init(common_params_speculative & params,
case COMMON_SPECULATIVE_TYPE_NGRAM_MAP_K:
case COMMON_SPECULATIVE_TYPE_NGRAM_MAP_K4V: {
impls.push_back(
std::make_unique<common_speculative_state_ngram_map_k>(
std::make_unique<common_speculative_impl_ngram_map_k>(
config.params, get_common_ngram_map(config.type, config.params.ngram_map_k), n_seq));
break;
}
case COMMON_SPECULATIVE_TYPE_NGRAM_MOD: {
impls.push_back(
std::make_unique<common_speculative_state_ngram_mod>(config.params, n_seq));
std::make_unique<common_speculative_impl_ngram_mod>(config.params, n_seq));
break;
}
case COMMON_SPECULATIVE_TYPE_NGRAM_CACHE: {
@ -975,7 +975,7 @@ common_speculative * common_speculative_init(common_params_speculative & params,
config, n_seq,
params.ngram_cache.lookup_cache_static,
params.ngram_cache.lookup_cache_dynamic);
impls.push_back(std::make_unique<common_speculative_state_ngram_cache>(state));
impls.push_back(std::make_unique<common_speculative_impl_ngram_cache>(state));
break;
}
default:
@ -984,7 +984,7 @@ common_speculative * common_speculative_init(common_params_speculative & params,
}
if (impls.empty()) {
LOG_WRN("%s", "no implementations specified for speculative decoding\n");
LOG_WRN("%s: no implementations specified for speculative decoding\n", __func__);
return nullptr;
}

14
convert_hf_to_gguf.py
View file

@ -2865,8 +2865,12 @@ class LlamaModel(TextModel):
# fix for SmolVLM2, missing `num_attention_heads` in config.json
if self.hf_arch == "VLlama3ForCausalLM":
self.hparams["num_attention_heads"] = self.hparams.get("num_attention_heads", 32)
hparams = ModelBase.load_hparams(self.dir_model, is_mistral_format=False)
self.origin_hf_arch = hparams.get('architectures', [None])[0]
# Mistral consolidated format has no config.json; origin_hf_arch is HF-only.
if self.is_mistral_format:
self.origin_hf_arch = None
else:
hparams = ModelBase.load_hparams(self.dir_model, is_mistral_format=False)
self.origin_hf_arch = hparams.get('architectures', [None])[0]
def set_vocab(self):
if self.origin_hf_arch == "GlmasrModel":
@ -13409,7 +13413,7 @@ class PixtralModel(LlavaVisionModel):
self.gguf_writer.add_vision_use_silu(True)
# spatial_merge_size
if self.find_vparam(["mm_projector_id"]) == "patch_merge":
if self.find_vparam(["mm_projector_id"], optional=True) == "patch_merge":
self.gguf_writer.add_vision_spatial_merge_size(
self.find_vparam(["spatial_merge_size"])
)
@ -13417,8 +13421,12 @@ class PixtralModel(LlavaVisionModel):
def map_tensor_name(self, name: str, try_suffixes: Sequence[str] = (".weight", ".bias")) -> str:
if name == "vision_language_adapter.w_in.weight":
return "mm.1.weight"
elif name == "vision_language_adapter.w_in.bias":
return "mm.1.bias"
elif name == "vision_language_adapter.w_out.weight":
return "mm.2.weight"
elif name == "vision_language_adapter.w_out.bias":
return "mm.2.bias"
return super().map_tensor_name(name, try_suffixes)

11
examples/llama-eval/llama-eval.py
View file

@ -1,5 +1,4 @@
#!/usr/bin/env python3
# type: ignore
import argparse
import json
@ -100,6 +99,8 @@ D) {D}
class BaseDataset(ABC):
questions: List[Dict]
@abstractmethod
def get_question(self, index: int) -> Dict:
pass
@ -573,7 +574,7 @@ def normalize_number(s: str) -> Optional[int]:
class AimeDataset(BaseDataset):
def __init__(self, split: str = "train"):
self.split = split
self.questions: List[Dict] = []
self.questions = []
self._load_dataset()
def _load_dataset(self):
@ -618,7 +619,7 @@ class AimeDataset(BaseDataset):
class Aime2025Dataset(BaseDataset):
def __init__(self):
self.questions: List[Dict] = []
self.questions = []
self._load_dataset()
def _load_dataset(self):
@ -681,7 +682,7 @@ class Aime2025Dataset(BaseDataset):
class Gsm8kDataset(BaseDataset):
def __init__(self, split: str = "test"):
self.split = split
self.questions: List[Dict] = []
self.questions = []
self._load_dataset()
def _load_dataset(self):
@ -742,7 +743,7 @@ class GpqaDataset(BaseDataset):
def __init__(self, variant: str = "diamond", seed: int = 1234):
self.variant = variant
self.seed = seed
self.questions: List[Dict] = []
self.questions = []
self._load_dataset()
def _load_dataset(self):

32
ggml/src/ggml-cpu/cmake/FindSMTIME.cmake Normal file
View file

@ -0,0 +1,32 @@
include(CheckCSourceRuns)
if (CMAKE_SYSTEM_PROCESSOR MATCHES "^(riscv)" AND GGML_CPU_RISCV64_SPACEMIT)
set(SMT_MARCH_STR "-march=rv64gcv_zfh_zvfh_zba_zicbop")
if (CMAKE_C_COMPILER_ID STREQUAL "GNU" AND
CMAKE_C_COMPILER_VERSION VERSION_GREATER_EQUAL 15)
string(APPEND SMT_MARCH_STR "_xsmtvdotii")
endif()
set(CMAKE_REQUIRED_FLAGS "${SMT_MARCH_STR}")
check_c_source_compiles("int main() {__asm__ volatile(\"vmadot v2, v0, v1\");}" SPACEMIT_RISCV_COMPILER_SUPPORT_IME1)
check_c_source_compiles("int main() {__asm__ volatile(\"vmadot v2, v0, v1, i4\");}" SPACEMIT_RISCV_COMPILER_SUPPORT_VMADOT_S4)
check_c_source_compiles("int main() {__asm__ volatile(\"vmadot v2, v0, v1, i8\");}" SPACEMIT_RISCV_COMPILER_SUPPORT_VMADOT_S8)
check_c_source_compiles("int main() {__asm__ volatile(\"vfwmadot v2, v0, v1, fp16\");}" SPACEMIT_RISCV_COMPILER_SUPPORT_VFWMADOT_FP16)
check_c_source_compiles("int main() {__asm__ volatile(\"vmadot.hp v2, v0, v1, v0, 0, i4\");}" SPACEMIT_RISCV_COMPILER_SUPPORT_VFMADOT_S4)
check_c_source_compiles("int main() {__asm__ volatile(\"vmadot.hp v2, v0, v1, v0, 0, i8\");}" SPACEMIT_RISCV_COMPILER_SUPPORT_VFMADOT_S8)
check_c_source_compiles("int main() {__asm__ volatile(\"vmadot1 v2, v0, v1\");}" SPACEMIT_RISCV_COMPILER_SUPPORT_VMADOTN)
check_c_source_compiles("int main() {__asm__ volatile(\"vpack.vv v2, v0, v1, 2\");}" SPACEMIT_RISCV_COMPILER_SUPPORT_VPACK)
check_c_source_compiles("int main() {__asm__ volatile(\"vnspack.vv v2, v0, v1, 2\");}" SPACEMIT_RISCV_COMPILER_SUPPORT_VNPACK)
unset(CMAKE_REQUIRED_FLAGS)
list(APPEND RISCV64_SPACEMIT_IME_SPEC "")
if (SPACEMIT_RISCV_COMPILER_SUPPORT_IME1)
set(RISCV64_SPACEMIT_IME_SPEC "RISCV64_SPACEMIT_IME1")
endif()
if (SPACEMIT_RISCV_COMPILER_SUPPORT_VMADOT_S4 AND SPACEMIT_RISCV_COMPILER_SUPPORT_VPACK AND SPACEMIT_RISCV_COMPILER_SUPPORT_VNPACK)
list(APPEND RISCV64_SPACEMIT_IME_SPEC "RISCV64_SPACEMIT_IME2")
endif()
message("RISCV64_SPACEMIT_IME_SPEC: ${RISCV64_SPACEMIT_IME_SPEC}")
endif()

12
ggml/src/ggml-cpu/ggml-cpu.c
View file

@ -51,6 +51,10 @@
#include "llamafile/sgemm.h"
#endif
#ifdef GGML_USE_CPU_RISCV64_SPACEMIT
# include "spacemit/ime.h"
#endif
// Note: once we move threading into a separate C++ file
// will use std::hardware_destructive_interference_size instead of hardcoding it here
// and we'll use C++ attribute syntax.
@ -3860,7 +3864,11 @@ static thread_ret_t ggml_graph_compute_thread(void * data) {
const struct ggml_cgraph * cgraph = tp->cgraph;
const struct ggml_cplan * cplan = tp->cplan;
#ifdef GGML_USE_CPU_RISCV64_SPACEMIT
ggml_backend_cpu_riscv64_spacemit_set_numa_thread_affinity(state->ith);
#else
set_numa_thread_affinity(state->ith);
#endif
struct ggml_compute_params params = {
/*.ith =*/ state->ith,
@ -3917,6 +3925,10 @@ static thread_ret_t ggml_graph_compute_thread(void * data) {
ggml_barrier(state->threadpool);
#ifdef GGML_USE_CPU_RISCV64_SPACEMIT
ggml_backend_cpu_riscv64_spacemit_clear_numa_thread_affinity_threaded(state->ith);
#endif
return 0;
}

2161
ggml/src/ggml-cpu/spacemit/ime.cpp
View file

File diff suppressed because it is too large Load diff

8
ggml/src/ggml-cpu/spacemit/ime.h
View file

@ -8,6 +8,14 @@ extern "C" {
ggml_backend_buffer_type_t ggml_backend_cpu_riscv64_spacemit_buffer_type(void);
void ggml_backend_cpu_riscv64_spacemit_set_numa_thread_affinity(int thread_n);
void ggml_backend_cpu_riscv64_spacemit_clear_numa_thread_affinity_threaded(int thread_n);
void * ggml_backend_cpu_riscv64_spacemit_alloc_shared(size_t size, size_t alignment);
void ggml_backend_cpu_riscv64_spacemit_free_shared(void * ptr);
#ifdef __cplusplus
}
#endif

3235
ggml/src/ggml-cpu/spacemit/ime1_kernels.cpp
View file

File diff suppressed because it is too large Load diff

5768
ggml/src/ggml-cpu/spacemit/ime2_kernels.cpp Normal file
View file

File diff suppressed because it is too large Load diff

320
ggml/src/ggml-cpu/spacemit/ime_env.cpp Normal file
View file

@ -0,0 +1,320 @@
#include "ime_env.h"
#include "ggml-impl.h"
#include "spine_mem_pool.h"
#include <fcntl.h>
#include <unistd.h>
#include <algorithm>
#include <array>
#include <cctype>
#include <fstream>
#include <string>
#include <thread>
#include <unordered_map>
namespace ggml::cpu::riscv64_spacemit {
bool spine_core_info::get_spine_core_info(std::vector<spine_core_info> & result) {
static std::unordered_map<uint64_t, spine_core_arch_id> spine_march_mapping_ = {
{0x8000000058000001, spine_core_arch_id::core_arch_x60 },
{ 0x8000000041000001, spine_core_arch_id::core_arch_a60 },
{ 0x8000000058000002, spine_core_arch_id::core_arch_x100},
{ 0x8000000041000002, spine_core_arch_id::core_arch_a100},
};
result.clear();
std::ifstream file("/proc/cpuinfo");
std::string line;
std::vector<std::array<uint64_t, 2>> cpu_info_list;
uint64_t current_processor = spine_invalid_core_id;
uint64_t current_marchid = 0;
bool has_processor = false;
bool has_marchid = false;
if (!file.is_open()) {
return false;
}
while (std::getline(file, line)) {
if (line.substr(0, 9) == "processor") {
if (has_processor && has_marchid) {
cpu_info_list.push_back({ current_processor, current_marchid });
}
size_t colon_pos = line.find(':');
if (colon_pos != std::string::npos) {
current_processor = std::stoi(line.substr(colon_pos + 1));
has_processor = true;
}
has_marchid = false;
} else if (line.substr(0, 7) == "marchid") {
size_t colon_pos = line.find(':');
if (colon_pos != std::string::npos) {
std::string marchid_str = line.substr(colon_pos + 1);
marchid_str.erase(std::remove_if(marchid_str.begin(), marchid_str.end(), isspace), marchid_str.end());
current_marchid = std::stoull(marchid_str, nullptr, 16);
has_marchid = true;
}
}
}
if (has_processor && has_marchid) {
cpu_info_list.push_back({ current_processor, current_marchid });
}
if (has_processor && has_marchid) {
for (auto & cpu_info : cpu_info_list) {
if (cpu_info[0] != spine_invalid_core_id &&
spine_march_mapping_.find(cpu_info[1]) != spine_march_mapping_.end()) {
auto core_info = spine_core_info();
core_info.core_id = cpu_info[0];
core_info.arch_id = spine_core_arch_id(spine_march_mapping_[cpu_info[1]]);
result.push_back(core_info);
}
}
}
return has_processor && has_marchid;
}
namespace {
uint16_t hex_string_to_u16(const std::string & hex_str) {
try {
size_t pos = 0;
if (hex_str.substr(0, 2) == "0x" || hex_str.substr(0, 2) == "0X") {
pos = 2;
}
unsigned long result = std::stoul(hex_str.substr(pos), nullptr, 16);
if (result > std::numeric_limits<uint16_t>::max()) {
throw std::out_of_range("Converted value is out of range for uint16_t");
}
return static_cast<uint16_t>(result);
} catch (const std::invalid_argument & e) {
throw std::invalid_argument("Invalid hexadecimal string");
} catch (const std::out_of_range & e) {
throw;
}
}
const char * spine_mem_pool_backend_to_string(spine_mem_pool_backend backend) {
switch (backend) {
case spine_mem_pool_backend::none:
return "NONE";
case spine_mem_pool_backend::posix_memalign:
return "POSIX";
case spine_mem_pool_backend::transparent_hugepage:
return "HPAGE";
case spine_mem_pool_backend::hugetlb_1g:
return "HPAGE1GB";
}
return "unknown";
}
spine_mem_pool_backend parse_mem_backend(const char * mem_backend_str) {
if (mem_backend_str == nullptr || mem_backend_str[0] == '\0') {
return spine_mem_pool_backend::transparent_hugepage;
}
std::string value(mem_backend_str);
std::transform(value.begin(), value.end(), value.begin(),
[](unsigned char ch) { return static_cast<char>(std::tolower(ch)); });
if (value == "none") {
return spine_mem_pool_backend::none;
}
if (value == "posix") {
return spine_mem_pool_backend::posix_memalign;
}
if (value == "hpage") {
return spine_mem_pool_backend::transparent_hugepage;
}
if (value == "hpage1gb") {
return spine_mem_pool_backend::hugetlb_1g;
}
throw std::runtime_error("invalid SPACEMIT_MEM_BACKEND: " + value + ", expected NONE, POSIX, HPAGE or HPAGE1GB");
}
} // namespace
spine_env_info::spine_env_info() {
num_cores = static_cast<int>(std::thread::hardware_concurrency());
spine_core_info::get_spine_core_info(core_info_list);
// special for x60 K1
if (core_info_list.size() == 8 && core_info_list[0].arch_id == spine_core_arch_id::core_arch_x60) {
for (int i = 0; i < 4; i++) {
core_info_list[i].arch_id = spine_core_arch_id::core_arch_a60;
}
}
// special for qemu
if (core_info_list.size() == 0) {
char * spine_core_arch_str = getenv("SPACEMIT_CORE_ARCH");
if (spine_core_arch_str != nullptr) {
auto arch_id = hex_string_to_u16(spine_core_arch_str);
for (int i = 0; i < num_cores; i++) {
auto core_info = spine_core_info();
core_info.core_id = i;
core_info.arch_id = spine_core_arch_id{ arch_id };
core_info_list.push_back(core_info);
}
}
}
if (core_info_list.size() == 0) {
throw std::runtime_error(
"Failed to get SPACEMIT_CORE_ARCH from environment or failed to parse it from /proc/cpuinfo");
}
char * spine_perfer_core_arch_str = getenv("SPACEMIT_PERFER_CORE_ARCH");
if (spine_perfer_core_arch_str != nullptr && spine_perfer_core_arch_str != "") {
perfer_core_arch_id = spine_core_arch_id{ hex_string_to_u16(spine_perfer_core_arch_str) };
}
char * spine_perfer_core_id_str = getenv("SPACEMIT_PERFER_CORE_ID");
std::vector<int> perfer_core_id_vec;
if (spine_perfer_core_id_str != nullptr && spine_perfer_core_id_str != "") {
std::string perfer_core_id_str(spine_perfer_core_id_str);
size_t start = 0;
size_t end = 0;
while ((end = perfer_core_id_str.find(',', start)) != std::string::npos) {
std::string core_id_substr = perfer_core_id_str.substr(start, end - start);
perfer_core_id_vec.push_back(std::stoi(core_id_substr));
start = end + 1;
}
std::string core_id_substr = perfer_core_id_str.substr(start);
perfer_core_id_vec.push_back(std::stoi(core_id_substr));
}
perfer_core_ids.reserve(num_cores);
if (perfer_core_arch_id == spine_core_arch_id::core_arch_none) {
for (auto & core_info : core_info_list) {
auto core_arch_id = core_info.arch_id;
auto core_arch_head = (uint16_t) (core_arch_id) >> 12;
if (core_arch_head == 0xA) {
num_perfer_cores++;
perfer_core_arch_id = core_arch_id;
cpu_mask |= (1ULL << core_info.core_id);
perfer_core_ids.push_back(core_info.core_id);
}
}
} else {
for (auto & core_info : core_info_list) {
auto core_arch_id = core_info.arch_id;
if (core_arch_id == perfer_core_arch_id) {
num_perfer_cores++;
cpu_mask |= (1ULL << core_info.core_id);
auto core_arch_head = (uint16_t) (core_arch_id) >> 12;
if (core_arch_head == 0xA) {
perfer_core_ids.push_back(core_info.core_id);
}
}
}
if (num_perfer_cores == 0) {
GGML_ABORT("can not find core with arch id %x for SPACEMIT_PERFER_CORE_ARCH in core info list\n",
(uint16_t) perfer_core_arch_id);
}
}
if (perfer_core_id_vec.size() > 0) {
perfer_core_ids.clear();
cpu_mask = 0;
num_perfer_cores = 0;
for (int core_id : perfer_core_id_vec) {
if (core_id < 0 || core_id >= num_cores) {
GGML_ABORT("invalid core id in SPACEMIT_PERFER_CORE_ID: %d, should be between 0 and %d\n", core_id,
num_cores - 1);
}
auto core_info = core_info_list[core_id];
auto core_arch_id = core_info.arch_id;
if (core_arch_id == perfer_core_arch_id) {
cpu_mask |= (1ULL << core_id);
perfer_core_ids.push_back(core_id);
} else {
GGML_ABORT(
"core id %d in SPACEMIT_PERFER_CORE_ID has arch id %x which does not match "
"SPACEMIT_PERFER_CORE_ARCH %x\n",
core_id, (uint16_t) core_arch_id, (uint16_t) perfer_core_arch_id);
}
}
std::string perfer_core_id_vec_str;
for (int core_id : perfer_core_id_vec) {
perfer_core_id_vec_str += std::to_string(core_id) + ",";
}
perfer_core_id_vec_str.pop_back();
GGML_LOG_DEBUG("SPACEMIT_PERFER_CORE_ID is set, perferred core ids: %s\n", perfer_core_id_vec_str.c_str());
num_perfer_cores = static_cast<int>(perfer_core_id_vec.size());
}
use_ime1 = perfer_core_arch_id == spine_core_arch_id::core_arch_a60 ||
perfer_core_arch_id == spine_core_arch_id::core_arch_x100;
use_ime2 = perfer_core_arch_id == spine_core_arch_id::core_arch_a100;
mem_backend = parse_mem_backend(getenv("SPACEMIT_MEM_BACKEND"));
char * spine_disable_tcm_str = getenv("SPACEMIT_DISABLE_TCM");
auto user_disable_tcm = spine_disable_tcm_str != nullptr && strcmp(spine_disable_tcm_str, "0") != 0;
if (!user_disable_tcm) {
spine_mem_pool_tcm_info tcm_info;
if (spine_mem_pool_tcm_init(&tcm_info)) {
use_tcm = tcm_info.available;
tcm_blk_size = tcm_info.blk_size;
GGML_LOG_DEBUG("CPU_RISCV64_SPACEMIT: tcm is available, blk_size: %zu, blk_num: %zu, is_fake_tcm: %d\n",
tcm_info.blk_size, tcm_info.blk_num, tcm_info.is_fake_tcm);
for (auto & core_info : core_info_list) {
auto core_arch_head = (uint16_t) (core_info.arch_id) >> 12;
if (core_arch_head != 0xA) {
aicpu_id_offset++;
} else {
break;
}
}
}
}
GGML_LOG_DEBUG(
"CPU_RISCV64_SPACEMIT: num_cores: %d, num_perfer_cores: %d, perfer_core_arch_id: %x, exclude_main_thread: %d, "
"use_ime1: %d, use_ime2: %d, mem_backend: %s, cpu_mask: %lx, aicpu_id_offset: %d\n",
num_cores, num_perfer_cores, (uint16_t) perfer_core_arch_id, exclude_main_thread, use_ime1, use_ime2,
spine_mem_pool_backend_to_string(mem_backend), cpu_mask, aicpu_id_offset);
const size_t init_barrier_size = sizeof(spine_barrier_t) * spine_init_barrier_count;
init_barrier =
static_cast<spine_barrier_t *>(spine_mem_pool_shared_mem_alloc(init_barrier_size, alignof(spine_barrier_t)));
if (init_barrier != nullptr) {
init_barrier_is_shared_mem = true;
} else {
GGML_LOG_WARN("CPU_RISCV64_SPACEMIT: failed to allocate init_barrier from shared mem, falling back to heap\n",
__func__);
init_barrier = new spine_barrier_t[spine_init_barrier_count];
}
spine_barrier_init(init_barrier, spine_init_barrier_count, 2);
}
spine_env_info::~spine_env_info() {
if (init_barrier_is_shared_mem) {
spine_mem_pool_shared_mem_free(init_barrier);
} else {
delete[] init_barrier;
}
init_barrier = nullptr;
init_barrier_is_shared_mem = false;
}
spine_env_info global_spine_env_info;
} // namespace ggml::cpu::riscv64_spacemit

55
ggml/src/ggml-cpu/spacemit/ime_env.h Normal file
View file

@ -0,0 +1,55 @@
#pragma once
#include "spine_barrier.h"
#include "spine_mem_pool.h"
#include <cstddef>
#include <cstdint>
#include <vector>
namespace ggml::cpu::riscv64_spacemit {
constexpr uint64_t spine_invalid_core_id = 0xFFFFFFFF;
constexpr size_t spine_init_barrier_count = 16;
enum class spine_core_arch_id : uint16_t {
core_arch_none = 0,
core_arch_x60 = 0x503C,
core_arch_x100 = 0x5064,
core_arch_x200 = 0x50C8,
core_arch_a60 = 0xA03C,
core_arch_a100 = 0xA064,
core_arch_a200 = 0xA0C8,
};
struct spine_core_info {
uint64_t core_id{ spine_invalid_core_id };
spine_core_arch_id arch_id{ spine_core_arch_id::core_arch_none };
static bool get_spine_core_info(std::vector<spine_core_info> & result);
};
struct spine_env_info {
std::vector<spine_core_info> core_info_list;
std::vector<int> perfer_core_ids;
int aicpu_id_offset{ 0 };
int num_cores{ 0 };
int num_perfer_cores{ 0 };
spine_core_arch_id perfer_core_arch_id{ spine_core_arch_id::core_arch_none };
bool exclude_main_thread{ false };
bool use_ime2{ false };
bool use_ime1{ false };
bool use_tcm{ false };
spine_mem_pool_backend mem_backend{ spine_mem_pool_backend::transparent_hugepage };
uint64_t tcm_blk_size{ 0 };
uint64_t cpu_mask{ 0 };
spine_barrier_t * init_barrier{ nullptr };
bool init_barrier_is_shared_mem{ false };
spine_env_info();
~spine_env_info();
};
extern spine_env_info global_spine_env_info;
} // namespace ggml::cpu::riscv64_spacemit

197
ggml/src/ggml-cpu/spacemit/ime_kernels.h
View file

@ -1,26 +1,189 @@
#pragma once
#include <cassert>
#include <cstddef>
#include <functional>
namespace spacemit_kernels {
#define BLOCK_QNK_LEN 256
template <int N> struct nrow_block_q2_k {
// [4bit scale + 4bit zp] * N * 16
uint8_t scales[N * BLOCK_QNK_LEN / 16];
// [b0, b16, b32, b48] [b1, b17, b33, b49] ... [b15, b31, b47, b63]
// [b64, b80, b96, b112] ...[b79, b95, b111, b127]
// [b128, b144, b160, b176] ...[b143, b159, b175, b191]
// [b192, b208, b224, b240] ...[b207, b223, b239, b255]
uint8_t qs[N * BLOCK_QNK_LEN / 4];
uint16_t scales16[N];
uint16_t zeros16[N];
};
template <int N> struct nrow_block_q3_k {
// [8bit scale] * N * 16
int8_t scales[N * 16];
// [b0, b1, b2, b3, b4, b5, b6, b7] ... [b248, b249, b250, b251, b252, b253, b254, b255]
uint8_t hmask[N * BLOCK_QNK_LEN / 8];
// [b0, b16, b32, b48] [b1, b17, b33, b49] ... [b15, b31, b47, b63]
// [b64, b80, b96, b112] ...[b79, b95, b111, b127]
// [b128, b144, b160, b176] ...[b143, b159, b175, b191]
// [b192, b208, b224, b240] ...[b207, b223, b239, b255]
uint8_t qs[N * BLOCK_QNK_LEN / 4];
uint16_t scales16[N];
};
template <int N> struct nrow_block_mxfp4 {
uint8_t e[N];
uint8_t qh[4 * N];
uint8_t qs[16 * N];
};
template <int N> struct __attribute__((packed)) nrow_block_q5_1 {
uint16_t scales16[N];
uint8_t zp[N];
// n0 [bh0, bh1, bh2, bh3, bh4, bh5, bh6, bh7] ....
uint8_t qh[4 * N];
// n0 [b0, b1], [b2, b3] .... [b30, b31]
// n1 [b0, b1], [b2, b3] .... [b30, b31]
uint8_t qs[16 * N];
};
static_assert(sizeof(nrow_block_q5_1<1>) == sizeof(uint8_t) + 22, "wrong nrow_block_q5_1 block size/padding");
template <int N> struct __attribute__((packed)) nrow_block_q5_0 {
uint16_t scales16[N];
// n0 [bh0, bh1, bh2, bh3, bh4, bh5, bh6, bh7] ....
uint8_t qh[4 * N];
// n0 [b0, b1], [b2, b3] .... [b30, b31]
// n1 [b0, b1], [b2, b3] .... [b30, b31]
uint8_t qs[16 * N];
};
static_assert(sizeof(nrow_block_q5_0<1>) == 22, "wrong nrow_block_q5_0 block size/padding");
using gemm_kernel_quantize_def = std::function<
size_t(size_t, const uint8_t *, const uint8_t *, const uint8_t *, float *, size_t, size_t, size_t, size_t)>;
using moe_gemm_kernel_quantize_def = std::function<
size_t(size_t, const uint8_t **, const uint8_t *, const uint8_t *, float **, size_t, size_t, size_t, size_t)>;
namespace sqnbitgemm_spacemit_ime {
namespace ime1 {
size_t gemm_kernel_i8i4(size_t blk_len,
const std::byte * quant_a_ptr,
const std::byte * quant_b_data,
const float * quant_b_scale,
const std::byte * quant_b_zp,
float * c_ptr,
size_t count_m,
size_t count_n,
size_t count_k,
size_t block_count_k,
size_t ldc,
const float * bias,
const size_t scale_stride);
size_t gemm_kernel_i8i4(size_t blk_len,
const uint8_t * quant_a_ptr,
const uint8_t * quant_b_data,
const uint8_t * quant_b_zp,
float * c_ptr,
size_t count_m,
size_t count_n,
size_t k_blks,
size_t ldc);
void quantize_a_row_i8(size_t blk_len, const float * a_ptr, size_t count_k, std::byte * quant_a_ptr);
void quantize_a_row_i8(size_t blk_len, const float * a_ptr, size_t count_k, uint8_t * quant_a_ptr);
void quantize_a_4row_i8(size_t blk_len, const float * a_ptr, size_t count_k, std::byte * quant_a_ptr);
void quantize_a_4row_i8(size_t blk_len, const float * a_ptr, size_t count_k, uint8_t * quant_a_ptr);
} // namespace ime1
} // namespace sqnbitgemm_spacemit_ime
namespace ime2 {
size_t gemm_kernel_i8i2k(size_t blk_len,
const uint8_t * quant_a_ptr,
const uint8_t * quant_b_data,
const uint8_t * quant_b_zp,
float * c_ptr,
size_t count_m,
size_t count_n,
size_t k_blks,
size_t ldc);
size_t gemm_kernel_i8i3k(size_t blk_len,
const uint8_t * quant_a_ptr,
const uint8_t * quant_b_data,
const uint8_t * quant_b_zp,
float * c_ptr,
size_t count_m,
size_t count_n,
size_t k_blks,
size_t ldc);
size_t gemm_kernel_i8i4(size_t blk_len,
const uint8_t * quant_a_ptr,
const uint8_t * quant_b_data,
const uint8_t * quant_b_zp,
float * c_ptr,
size_t count_m,
size_t count_n,
size_t k_blks,
size_t ldc);
size_t gemm_kernel_i8i4_hp(size_t blk_len,
const uint8_t * quant_a_ptr,
const uint8_t * quant_b_data,
const uint8_t * quant_b_zp,
float * c_ptr,
size_t count_m,
size_t count_n,
size_t k_blks,
size_t ldc);
size_t moe_m2_gemm_kernel_i8i4(size_t blk_len,
const uint8_t ** quant_a_ptr,
const uint8_t * quant_b_data,
const uint8_t * quant_b_zp,
float ** c_ptr,
size_t count_m,
size_t count_n,
size_t k_blks,
size_t ldc);
size_t gemm_kernel_i8i8(size_t blk_len,
const uint8_t * quant_a_ptr,
const uint8_t * quant_b_data,
const uint8_t * quant_b_zp,
float * c_ptr,
size_t count_m,
size_t count_n,
size_t k_blks,
size_t ldc);
size_t gemm_kernel_i8mxfp4(size_t blk_len,
const uint8_t * quant_a_ptr,
const uint8_t * quant_b_data,
const uint8_t * quant_b_zp,
float * c_ptr,
size_t count_m,
size_t count_n,
size_t k_blks,
size_t ldc);
size_t moe_m2_gemm_kernel_i8mxfp4(size_t blk_len,
const uint8_t ** quant_a_ptr,
const uint8_t * quant_b_data,
const uint8_t * quant_b_zp,
float ** c_ptr,
size_t count_m,
size_t count_n,
size_t k_blks,
size_t ldc);
size_t gemm_kernel_i8i5(size_t blk_len,
const uint8_t * quant_a_ptr,
const uint8_t * quant_b_data,
const uint8_t * quant_b_zp,
float * c_ptr,
size_t count_m,
size_t count_n,
size_t k_blks,
size_t ldc);
size_t moe_m2_gemm_kernel_i8i5(size_t blk_len,
const uint8_t ** quant_a_ptr,
const uint8_t * quant_b_data,
const uint8_t * quant_b_zp,
float ** c_ptr,
size_t count_m,
size_t count_n,
size_t k_blks,
size_t ldc);
} // namespace ime2
} // namespace spacemit_kernels

1795
ggml/src/ggml-cpu/spacemit/repack.cpp Normal file
View file

File diff suppressed because it is too large Load diff

14
ggml/src/ggml-cpu/spacemit/repack.h Normal file
View file

@ -0,0 +1,14 @@
#pragma once
#include "ggml-common.h"
#include "ggml.h"
#include <cstddef>
#include <cstdint>
namespace ggml::cpu::riscv64_spacemit {
template <typename BLOC_TYPE, int64_t INTER_SIZE, int64_t NB_COLS>
int repack(ggml_tensor * t, const void * data, size_t data_size);
} // namespace ggml::cpu::riscv64_spacemit

3178
ggml/src/ggml-cpu/spacemit/rvv_kernels.cpp Normal file
View file

File diff suppressed because it is too large Load diff

95
ggml/src/ggml-cpu/spacemit/rvv_kernels.h Normal file
View file

@ -0,0 +1,95 @@
#pragma once
#include "ggml-cpu-impl.h"
#include <cassert>
#include <cstddef>
#include <cstdint>
#include <functional>
namespace spacemit_kernels {
constexpr auto div_round_up(auto up, auto down) {
return (up + down - 1) / down;
}
// Q8 Blk [f32] [s16] [int8 * blk_len]
// Q8 Blk N [f32 * N] [s16 * N] [int8 * blk_len * N]
constexpr size_t q8_blk_size(size_t blk_len, bool with_blk_sum = false) {
const size_t blk_size = sizeof(float) + blk_len * sizeof(int8_t) + (with_blk_sum ? sizeof(int16_t) : 0);
return blk_size;
}
// Q8 HP row block: K is split into K32 subblocks.
// Each subblock stores [f32 scale] [int8 * 32], with an optional fp16 sum trailer per subblock.
constexpr size_t q8_hp_blk_size(size_t blk_len, bool with_blk_sum = false, bool with_blk_scale = false) {
const size_t subblk_count = div_round_up(blk_len, size_t(32));
const size_t blk_size = blk_len * sizeof(int8_t) + subblk_count * sizeof(_Float16) +
(with_blk_sum ? subblk_count * sizeof(_Float16) : 0) +
(with_blk_scale ? sizeof(_Float16) : 0);
return blk_size;
}
// Q8K Blk [f32] [s16 * (blk_len / 16)] [int8 * blk_len]
// Q8K Blk N [f32 * N] [s16 * (blk_len / 16) * N] [int8 * blk_len * N]
constexpr size_t q8k_blk_size(size_t blk_len) {
const size_t blk_size = sizeof(float) + blk_len * sizeof(int8_t) + sizeof(int16_t) * blk_len / 16;
return blk_size;
}
using quantize_a_row_def = std::function<void(size_t, const float *, size_t, uint8_t *)>;
namespace rvv {
void memcpy1d(void * dst, const void * src, int64_t size);
void memcpy2d(void * dst, int64_t dst_stride, const void * src, int64_t src_stride, int64_t tile_rows, int64_t size);
void forward_flash_attn_ext_f16_one_chunk_vlen1024_vf16(const ggml_compute_params * params,
ggml_tensor * dst,
int ir0,
int ir1,
void * tcm_buffer,
size_t tcm_buffer_size);
void forward_flash_attn_ext_f16_tiled_vlen1024_vf16(const ggml_compute_params * params,
ggml_tensor * dst,
int ir0,
int ir1,
void * tcm_buffer,
size_t tcm_buffer_size);
void forward_rms_norm_f32(ggml_compute_params * params, ggml_tensor * op);
void forward_norm_f32(ggml_compute_params * params, ggml_tensor * op);
void forward_cont_with_permute(ggml_compute_params * params, ggml_tensor * op);
void forward_cpy_with_permute(ggml_compute_params * params, ggml_tensor * op);
template <typename T> void forward_get_rows(ggml_compute_params * params, ggml_tensor * op);
template <typename T> void forward_concat(ggml_compute_params * params, ggml_tensor * op);
template <ggml_op op_type, typename T> void forward_binary(ggml_compute_params * params, ggml_tensor * op);
template <typename T> void forward_sum_rows(const ggml_compute_params * params, ggml_tensor * op);
template <typename T> void forward_repeat_nrows(ggml_compute_params * params, ggml_tensor * op);
template <typename T> void forward_repeat_dim1(ggml_compute_params * params, ggml_tensor * op);
void quantize_a_row_i8(size_t blk_len, const float * a_ptr, size_t count_k, uint8_t * quant_a_ptr);
void quantize_a_4row_i8(size_t blk_len, const float * a_ptr, size_t count_k, uint8_t * quant_a_ptr);
void quantize_a_row_i8_hp(size_t blk_len, const float * a_ptr, size_t count_k, uint8_t * quant_a_ptr);
void quantize_a_4row_i8_hp(size_t blk_len, const float * a_ptr, size_t count_k, uint8_t * quant_a_ptr);
void quantize_a_row_i8k(size_t blk_len, const float * a_ptr, size_t count_k, uint8_t * quant_a_ptr);
void quantize_a_4row_i8k(size_t blk_len, const float * a_ptr, size_t count_k, uint8_t * quant_a_ptr);
} // namespace rvv
} // namespace spacemit_kernels

34
ggml/src/ggml-cpu/spacemit/spine_barrier.h Normal file
View file

@ -0,0 +1,34 @@
#pragma once
#include <atomic>
#include <cstdint>
#define SPINE_CACHE_LINE 64
#define SPINE_CACHE_ALIGN __attribute__((aligned(SPINE_CACHE_LINE)))
struct spine_barrier_t {
SPINE_CACHE_ALIGN std::atomic<int64_t> pending_;
SPINE_CACHE_ALIGN std::atomic<int64_t> rounds_;
SPINE_CACHE_ALIGN int64_t total_;
};
inline void spine_barrier_wait(spine_barrier_t * b) {
auto cur_round = b->rounds_.load(std::memory_order_acquire);
auto cnt = --b->pending_;
if (cnt == 0) {
b->pending_.store(b->total_);
b->rounds_.store(cur_round + 1);
} else {
while (cur_round == b->rounds_.load(std::memory_order_relaxed)) {
__asm__ volatile("pause " ::: "memory");
}
}
}
inline void spine_barrier_init(spine_barrier_t * b, int num_barriers, uint64_t thread_count) {
for (int i = 0; i < num_barriers; i++) {
b[i].total_ = thread_count;
b[i].pending_.store(thread_count);
b[i].rounds_.store(0);
}
}

760
ggml/src/ggml-cpu/spacemit/spine_mem_pool.cpp Normal file
View file

@ -0,0 +1,760 @@
#include "spine_mem_pool.h"
#include "common.h"
#include "ime_env.h"
#include "spine_tcm.h"
#include <fcntl.h>
#include <sys/ioctl.h>
#include <sys/mman.h>
#include <unistd.h>
#include <algorithm>
#include <cerrno>
#include <cstdint>
#include <cstdlib>
#include <limits>
#include <memory>
#include <mutex>
#include <unordered_map>
#include <vector>
namespace ggml::cpu::riscv64_spacemit {
namespace {
constexpr size_t SPINE_MEM_POOL_CHUNK_SIZE = 512ull * 1024ull * 1024ull;
constexpr size_t SPINE_SHARE_MEM_POOL_CHUNK_SIZE = 512ull * 1024ull;
constexpr size_t SPINE_MEM_POOL_1G_REGION_SIZE = 1ull << 30;
constexpr uint64_t HUGETLB_1G_FLAG_REQUIRE_PUD = 1ull << 0;
constexpr char SPINE_MEM_POOL_HUGETLB_1G_DEV[] = "/dev/hugetlb_1g";
constexpr char SPINE_MEM_POOL_TCM_SYNC_MEM_DEV[] = "/dev/tcm_sync_mem";
struct hugetlb_1g_region {
uint64_t size{ 0 };
uint64_t dma_addr{ 0 };
uint64_t flags{ 0 };
uint64_t reserved{ 0 };
};
#define HUGETLB_1G_IOC_MAGIC 'M'
#define HUGETLB_1G_IOC_ALLOC _IOWR(HUGETLB_1G_IOC_MAGIC, 0x00, struct hugetlb_1g_region)
#define HUGETLB_1G_IOC_FREE _IO(HUGETLB_1G_IOC_MAGIC, 0x01)
struct free_block {
size_t offset{ 0 };
size_t size{ 0 };
};
struct pool_chunk {
uint8_t * base{ nullptr };
size_t size{ 0 };
int fd{ -1 };
std::vector<free_block> free_blocks;
};
struct pool_allocation {
void * chunk_base{ nullptr };
size_t chunk_size{ 0 };
void * base{ nullptr };
size_t size{ 0 };
};
bool is_power_of_two(size_t value) {
return value != 0 && (value & (value - 1)) == 0;
}
bool align_up(size_t value, size_t alignment, size_t * aligned_value) {
if (aligned_value == nullptr || alignment == 0) {
return false;
}
const size_t remainder = value % alignment;
if (remainder == 0) {
*aligned_value = value;
return true;
}
const size_t padding = alignment - remainder;
if (value > std::numeric_limits<size_t>::max() - padding) {
return false;
}
*aligned_value = value + padding;
return true;
}
bool align_up_uintptr(uintptr_t value, size_t alignment, uintptr_t * aligned_value) {
if (aligned_value == nullptr || alignment == 0) {
return false;
}
const uintptr_t remainder = value % alignment;
if (remainder == 0) {
*aligned_value = value;
return true;
}
const uintptr_t padding = alignment - remainder;
if (value > std::numeric_limits<uintptr_t>::max() - padding) {
return false;
}
*aligned_value = value + padding;
return true;
}
class spine_mem_pool_manager {
public:
explicit spine_mem_pool_manager(size_t default_chunk_size) : default_chunk_size_(default_chunk_size) {}
virtual ~spine_mem_pool_manager() = default;
void * alloc(size_t size, size_t alignment) {
if (size == 0 || !is_power_of_two(alignment)) {
return nullptr;
}
size_t aligned_size = 0;
if (!align_up(size, alignment, &aligned_size)) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: align_up failed for size %zu alignment %zu\n", __func__, size,
alignment);
return nullptr;
}
pool_allocation allocation;
std::lock_guard<std::mutex> lock(mutex_);
if (!try_alloc_locked(aligned_size, alignment, &allocation)) {
if (!add_chunk_locked(aligned_size, alignment)) {
return nullptr;
}
if (!try_alloc_locked(aligned_size, alignment, &allocation)) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: allocation retry failed for size %zu alignment %zu\n",
__func__, aligned_size, alignment);
return nullptr;
}
}
try {
const auto [allocation_it, inserted] = allocations_.emplace(allocation.base, allocation);
if (!inserted) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: duplicate allocation key %p\n", __func__, allocation.base);
rollback_allocation_locked(allocation);
return nullptr;
}
} catch (const std::bad_alloc &) {
rollback_allocation_locked(allocation);
throw;
}
return allocation.base;
}
void free(void * base) {
if (base == nullptr) {
return;
}
std::lock_guard<std::mutex> lock(mutex_);
auto allocation_it = allocations_.find(base);
if (allocation_it == allocations_.end()) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: unknown allocation %p\n", __func__, base);
return;
}
pool_allocation allocation = allocation_it->second;
allocations_.erase(allocation_it);
auto chunk_it = find_chunk_locked(allocation);
if (chunk_it == chunks_.end()) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: unknown chunk for allocation %p size %zu\n", __func__,
allocation.base, allocation.size);
return;
}
auto * chunk_base = chunk_it->base;
auto * alloc_base = static_cast<uint8_t *>(allocation.base);
if (alloc_base < chunk_base || alloc_base >= chunk_base + chunk_it->size) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: allocation %p out of chunk range %p..%p\n", __func__,
allocation.base, chunk_base, chunk_base + chunk_it->size);
return;
}
const size_t offset = static_cast<size_t>(alloc_base - chunk_base);
if (offset > chunk_it->size || allocation.size > chunk_it->size - offset) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: allocation %p size %zu exceeds chunk size %zu\n", __func__,
allocation.base, allocation.size, chunk_it->size);
return;
}
insert_free_block_locked(*chunk_it, { offset, allocation.size });
maybe_release_empty_chunk_locked(chunk_it);
}
protected:
void release_chunks() {
std::lock_guard<std::mutex> lock(mutex_);
allocations_.clear();
for (auto & chunk : chunks_) {
dealloc_chunk(&chunk);
}
chunks_.clear();
}
size_t default_chunk_size() const { return default_chunk_size_; }
static void clear_chunk(pool_chunk * chunk) {
chunk->base = nullptr;
chunk->size = 0;
chunk->fd = -1;
chunk->free_blocks.clear();
}
virtual bool alloc_chunk(size_t min_size, size_t alignment, void * hint_addr, pool_chunk * chunk) = 0;
virtual void dealloc_chunk(pool_chunk * chunk) = 0;
private:
struct alloc_candidate {
size_t chunk_index{ 0 };
size_t block_index{ 0 };
size_t aligned_offset{ 0 };
uintptr_t address{ std::numeric_limits<uintptr_t>::max() };
bool valid{ false };
};
std::vector<pool_chunk>::iterator find_chunk_locked(const pool_allocation & allocation) {
return std::find_if(chunks_.begin(), chunks_.end(), [&](const pool_chunk & chunk) {
return chunk.base == allocation.chunk_base && chunk.size == allocation.chunk_size;
});
}
bool add_chunk_locked(size_t min_size, size_t alignment) {
pool_chunk chunk;
const size_t chunk_request = default_chunk_size_ == 0 ? min_size : std::max(min_size, default_chunk_size_);
void * hint_addr = nullptr;
for (const auto & existing_chunk : chunks_) {
auto * chunk_end = existing_chunk.base + existing_chunk.size;
if (hint_addr == nullptr || chunk_end > hint_addr) {
hint_addr = chunk_end;
}
}
if (!alloc_chunk(chunk_request, alignment, hint_addr, &chunk)) {
return false;
}
if (chunk.base == nullptr || chunk.size < min_size) {
GGML_LOG_ERROR(
"CPU_RISCV64_SPACEMIT: %s: invalid chunk returned for request size %zu, chunk_base=%p chunk_size=%zu\n",
__func__, min_size, chunk.base, chunk.size);
dealloc_chunk(&chunk);
return false;
}
try {
chunk.free_blocks.push_back({ 0, chunk.size });
chunks_.push_back(std::move(chunk));
} catch (const std::bad_alloc &) {
dealloc_chunk(&chunk);
throw;
}
return true;
}
void rollback_allocation_locked(const pool_allocation & allocation) {
auto chunk_it = find_chunk_locked(allocation);
if (chunk_it == chunks_.end()) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: failed to rollback allocation %p, owning chunk not found\n",
__func__, allocation.base);
return;
}
auto * chunk_base = chunk_it->base;
auto * alloc_base = static_cast<uint8_t *>(allocation.base);
if (alloc_base < chunk_base || alloc_base >= chunk_base + chunk_it->size) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: failed to rollback allocation %p, chunk range is invalid\n",
__func__, allocation.base);
return;
}
const size_t offset = static_cast<size_t>(alloc_base - chunk_base);
if (offset > chunk_it->size || allocation.size > chunk_it->size - offset) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: failed to rollback allocation %p size %zu\n", __func__,
allocation.base, allocation.size);
return;
}
insert_free_block_locked(*chunk_it, { offset, allocation.size });
maybe_release_empty_chunk_locked(chunk_it);
}
bool try_alloc_locked(size_t size, size_t alignment, pool_allocation * allocation) {
alloc_candidate best;
for (size_t chunk_index = 0; chunk_index < chunks_.size(); ++chunk_index) {
const auto & chunk = chunks_[chunk_index];
for (size_t block_index = 0; block_index < chunk.free_blocks.size(); ++block_index) {
const auto & block = chunk.free_blocks[block_index];
uintptr_t aligned_addr = 0;
const auto block_addr = reinterpret_cast<uintptr_t>(chunk.base + block.offset);
if (!align_up_uintptr(block_addr, alignment, &aligned_addr)) {
continue;
}
if (aligned_addr < block_addr) {
continue;
}
const size_t aligned_offset = block.offset + static_cast<size_t>(aligned_addr - block_addr);
const size_t padding = aligned_offset - block.offset;
if (padding > block.size || size > block.size - padding) {
continue;
}
if (!best.valid || aligned_addr < best.address) {
best.chunk_index = chunk_index;
best.block_index = block_index;
best.aligned_offset = aligned_offset;
best.address = aligned_addr;
best.valid = true;
}
}
}
if (!best.valid) {
return false;
}
auto & chunk = chunks_[best.chunk_index];
const free_block block = chunk.free_blocks[best.block_index];
const size_t padding = best.aligned_offset - block.offset;
const size_t alloc_end = best.aligned_offset + size;
const size_t block_end = block.offset + block.size;
chunk.free_blocks.erase(chunk.free_blocks.begin() + best.block_index);
auto insert_it = chunk.free_blocks.begin() + best.block_index;
if (padding != 0) {
insert_it = chunk.free_blocks.insert(insert_it, { block.offset, padding });
++insert_it;
}
if (alloc_end < block_end) {
chunk.free_blocks.insert(insert_it, { alloc_end, block_end - alloc_end });
}
allocation->chunk_base = chunk.base;
allocation->chunk_size = chunk.size;
allocation->base = chunk.base + best.aligned_offset;
allocation->size = size;
return true;
}
void maybe_release_empty_chunk_locked(std::vector<pool_chunk>::iterator chunk_it) {
if (chunk_it->free_blocks.size() != 1) {
return;
}
const auto & block = chunk_it->free_blocks.front();
if (block.offset != 0 || block.size != chunk_it->size) {
return;
}
dealloc_chunk(&*chunk_it);
chunks_.erase(chunk_it);
}
void insert_free_block_locked(pool_chunk & chunk, free_block block) {
auto it = chunk.free_blocks.begin();
while (it != chunk.free_blocks.end() && it->offset < block.offset) {
++it;
}
if (it != chunk.free_blocks.begin()) {
const auto & prev = *(it - 1);
if (prev.offset + prev.size > block.offset) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: overlapping free block at offset %zu size %zu\n", __func__,
block.offset, block.size);
return;
}
}
if (it != chunk.free_blocks.end() && block.offset + block.size > it->offset) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: overlapping next free block at offset %zu size %zu\n", __func__,
block.offset, block.size);
return;
}
it = chunk.free_blocks.insert(it, block);
if (it != chunk.free_blocks.begin()) {
auto prev = it - 1;
if (prev->offset + prev->size == it->offset) {
it->offset = prev->offset;
it->size += prev->size;
it = chunk.free_blocks.erase(prev);
}
}
if (it + 1 != chunk.free_blocks.end() && it->offset + it->size == (it + 1)->offset) {
it->size += (it + 1)->size;
chunk.free_blocks.erase(it + 1);
}
}
std::mutex mutex_;
std::vector<pool_chunk> chunks_;
std::unordered_map<void *, pool_allocation> allocations_;
size_t default_chunk_size_{ 0 };
};
class spine_mem_pool_posix final : public spine_mem_pool_manager {
public:
spine_mem_pool_posix() : spine_mem_pool_manager(0) {}
~spine_mem_pool_posix() override { release_chunks(); }
private:
bool alloc_chunk(size_t min_size, size_t alignment, void * hint_addr, pool_chunk * chunk) override {
(void) hint_addr;
const size_t alloc_alignment = std::max(alignment, sizeof(void *));
void * base = nullptr;
const int rc = posix_memalign(&base, alloc_alignment, min_size);
if (rc != 0) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: posix_memalign failed for size %zu alignment %zu, rc=%d\n",
__func__, min_size, alloc_alignment, rc);
return false;
}
chunk->base = static_cast<uint8_t *>(base);
chunk->size = min_size;
chunk->fd = -1;
return true;
}
void dealloc_chunk(pool_chunk * chunk) override {
std::free(chunk->base);
clear_chunk(chunk);
}
};
class spine_mem_pool_transparent_hugepage final : public spine_mem_pool_manager {
public:
spine_mem_pool_transparent_hugepage() : spine_mem_pool_manager(SPINE_MEM_POOL_CHUNK_SIZE) {}
~spine_mem_pool_transparent_hugepage() override { release_chunks(); }
private:
bool alloc_chunk(size_t min_size, size_t alignment, void * hint_addr, pool_chunk * chunk) override {
(void) alignment;
size_t chunk_size = 0;
if (!align_up(min_size, default_chunk_size(), &chunk_size)) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: failed to round chunk size for %zu\n", __func__, min_size);
return false;
}
void * map_addr = mmap(hint_addr, chunk_size, PROT_READ | PROT_WRITE, MAP_PRIVATE | MAP_ANONYMOUS, -1, 0);
if (map_addr == MAP_FAILED) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: mmap failed for chunk size %zu, errno=%d\n", __func__, chunk_size,
errno);
return false;
}
if (madvise(map_addr, chunk_size, MADV_HUGEPAGE) != 0) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: madvise(MADV_HUGEPAGE) failed for chunk size %zu, errno=%d\n",
__func__, chunk_size, errno);
munmap(map_addr, chunk_size);
return false;
}
chunk->base = static_cast<uint8_t *>(map_addr);
chunk->size = chunk_size;
chunk->fd = -1;
return true;
}
void dealloc_chunk(pool_chunk * chunk) override {
if (chunk->base != nullptr && chunk->size != 0 && munmap(chunk->base, chunk->size) != 0) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: munmap failed for chunk %p size %zu, errno=%d\n", __func__,
chunk->base, chunk->size, errno);
}
clear_chunk(chunk);
}
};
class spine_mem_pool_hugetlb_1g final : public spine_mem_pool_manager {
public:
spine_mem_pool_hugetlb_1g() : spine_mem_pool_manager(SPINE_MEM_POOL_1G_REGION_SIZE) {}
~spine_mem_pool_hugetlb_1g() override { release_chunks(); }
private:
bool alloc_chunk(size_t min_size, size_t alignment, void * hint_addr, pool_chunk * chunk) override {
(void) alignment;
(void) hint_addr;
size_t region_size = 0;
if (!align_up(min_size, SPINE_MEM_POOL_1G_REGION_SIZE, &region_size)) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: failed to round hugetlb_1g size for %zu\n", __func__, min_size);
return false;
}
const int fd = open(SPINE_MEM_POOL_HUGETLB_1G_DEV, O_RDWR);
if (fd < 0) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: open(%s) failed, errno=%d\n", __func__,
SPINE_MEM_POOL_HUGETLB_1G_DEV, errno);
return false;
}
hugetlb_1g_region region;
region.size = region_size;
region.flags = HUGETLB_1G_FLAG_REQUIRE_PUD;
if (ioctl(fd, HUGETLB_1G_IOC_ALLOC, &region) < 0) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: HUGETLB_1G_IOC_ALLOC failed for size %zu, errno=%d\n", __func__,
region_size, errno);
close(fd);
return false;
}
void * map_addr = mmap(nullptr, region.size, PROT_READ | PROT_WRITE, MAP_SHARED, fd, 0);
if (map_addr == MAP_FAILED) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: mmap failed for hugetlb_1g size %llu, errno=%d\n", __func__,
static_cast<unsigned long long>(region.size), errno);
ioctl(fd, HUGETLB_1G_IOC_FREE);
close(fd);
return false;
}
chunk->base = static_cast<uint8_t *>(map_addr);
chunk->size = region.size;
chunk->fd = fd;
return true;
}
void dealloc_chunk(pool_chunk * chunk) override {
if (chunk->base != nullptr && chunk->size != 0 && munmap(chunk->base, chunk->size) != 0) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: munmap failed for hugetlb_1g chunk %p size %zu, errno=%d\n",
__func__, chunk->base, chunk->size, errno);
}
if (chunk->fd >= 0) {
if (ioctl(chunk->fd, HUGETLB_1G_IOC_FREE) < 0) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: HUGETLB_1G_IOC_FREE failed for chunk %p, errno=%d\n",
__func__, chunk->base, errno);
}
close(chunk->fd);
}
clear_chunk(chunk);
}
};
class spine_mem_pool_shared_mem final : public spine_mem_pool_manager {
public:
spine_mem_pool_shared_mem() : spine_mem_pool_manager(SPINE_SHARE_MEM_POOL_CHUNK_SIZE) {}
~spine_mem_pool_shared_mem() override { release_chunks(); }
private:
bool alloc_chunk(size_t min_size, size_t alignment, void * hint_addr, pool_chunk * chunk) override {
(void) alignment;
if (hint_addr != nullptr) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: shared_mem does not support multiple active chunks\n", __func__);
return false;
}
if (min_size > default_chunk_size()) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: shared_mem request %zu exceeds chunk size %zu\n", __func__,
min_size, default_chunk_size());
return false;
}
const int fd = open(SPINE_MEM_POOL_TCM_SYNC_MEM_DEV, O_RDWR | O_SYNC);
if (fd < 0) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: open(%s) failed, errno=%d\n", __func__,
SPINE_MEM_POOL_TCM_SYNC_MEM_DEV, errno);
return false;
}
void * map_addr = mmap(nullptr, default_chunk_size(), PROT_READ | PROT_WRITE, MAP_SHARED, fd, 0);
if (map_addr == MAP_FAILED) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: mmap failed for %s size %zu, errno=%d\n", __func__,
SPINE_MEM_POOL_TCM_SYNC_MEM_DEV, default_chunk_size(), errno);
close(fd);
return false;
}
chunk->base = static_cast<uint8_t *>(map_addr);
chunk->size = default_chunk_size();
chunk->fd = fd;
return true;
}
void dealloc_chunk(pool_chunk * chunk) override {
if (chunk->base != nullptr && chunk->size != 0 && munmap(chunk->base, chunk->size) != 0) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: munmap failed for shared_mem chunk %p size %zu, errno=%d\n",
__func__, chunk->base, chunk->size, errno);
}
if (chunk->fd >= 0) {
close(chunk->fd);
}
clear_chunk(chunk);
}
};
spine_mem_pool_manager & get_spine_mem_pool_manager() {
static std::once_flag pool_once;
static std::unique_ptr<spine_mem_pool_manager> selected_pool;
static spine_mem_pool_backend selected_backend = spine_mem_pool_backend::none;
spine_mem_pool_backend backend = global_spine_env_info.mem_backend;
if (backend == spine_mem_pool_backend::none) {
backend = spine_mem_pool_backend::transparent_hugepage;
}
std::call_once(pool_once, [&]() {
selected_backend = backend;
switch (selected_backend) {
case spine_mem_pool_backend::posix_memalign:
selected_pool = std::make_unique<spine_mem_pool_posix>();
break;
case spine_mem_pool_backend::transparent_hugepage:
selected_pool = std::make_unique<spine_mem_pool_transparent_hugepage>();
break;
case spine_mem_pool_backend::hugetlb_1g:
selected_pool = std::make_unique<spine_mem_pool_hugetlb_1g>();
break;
case spine_mem_pool_backend::none:
selected_backend = spine_mem_pool_backend::transparent_hugepage;
selected_pool = std::make_unique<spine_mem_pool_transparent_hugepage>();
break;
}
});
if (backend != selected_backend) {
GGML_LOG_ERROR(
"CPU_RISCV64_SPACEMIT: %s: mem pool backend is process-global and mutually exclusive, requested=%d but "
"selected=%d\n",
__func__, static_cast<int>(backend), static_cast<int>(selected_backend));
}
if (selected_pool) {
return *selected_pool;
}
throw std::bad_alloc();
}
spine_mem_pool_manager & get_spine_mem_pool_shared_mem_manager() {
static std::once_flag shared_mem_pool_once;
static std::unique_ptr<spine_mem_pool_shared_mem> shared_mem_pool;
std::call_once(shared_mem_pool_once, [&]() { shared_mem_pool = std::make_unique<spine_mem_pool_shared_mem>(); });
if (shared_mem_pool) {
return *shared_mem_pool;
}
throw std::bad_alloc();
}
} // namespace
bool spine_mem_pool_tcm_init(spine_mem_pool_tcm_info * info) noexcept {
if (info == nullptr) {
return false;
}
*info = {};
if (spine_tcm_open_handle(NULL) != 0 || !spine_tcm_is_available()) {
return false;
}
spine_tcm_mem_info_t mem_info;
if (spine_tcm_mem_info(&mem_info) != 0) {
return false;
}
info->available = true;
info->blk_size = mem_info.blk_size;
info->blk_num = mem_info.blk_num;
info->is_fake_tcm = mem_info.is_fake_tcm != 0;
return true;
}
void * spine_mem_pool_tcm_mem_get(int cpu_id) noexcept {
return spine_tcm_mem_get(cpu_id);
}
void * spine_mem_pool_tcm_mem_wait(int cpu_id) noexcept {
return spine_tcm_mem_try_wait(cpu_id, 1000 * 1000);
}
int spine_mem_pool_tcm_mem_release(int cpu_id) noexcept {
return spine_tcm_mem_release(cpu_id);
}
void * spine_mem_pool_alloc(size_t size, size_t alignment) noexcept {
try {
return get_spine_mem_pool_manager().alloc(size, alignment);
} catch (const std::bad_alloc &) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: bad_alloc while allocating size %zu\n", __func__, size);
return nullptr;
}
}
void * spine_mem_pool_shared_mem_alloc(size_t size, size_t alignment) noexcept {
try {
return get_spine_mem_pool_shared_mem_manager().alloc(size, alignment);
} catch (const std::bad_alloc &) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: bad_alloc while allocating shared memory size %zu\n", __func__, size);
return nullptr;
}
}
void spine_mem_pool_free(void * base) noexcept {
try {
get_spine_mem_pool_manager().free(base);
} catch (const std::bad_alloc &) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: bad_alloc while freeing allocation %p\n", __func__, base);
}
}
void spine_mem_pool_shared_mem_free(void * base) noexcept {
try {
get_spine_mem_pool_shared_mem_manager().free(base);
} catch (const std::bad_alloc &) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: bad_alloc while freeing shared allocation %p\n", __func__, base);
}
}
} // namespace ggml::cpu::riscv64_spacemit
extern "C" {
void * ggml_backend_cpu_riscv64_spacemit_alloc_shared(size_t size, size_t alignment) {
void * result = ggml::cpu::riscv64_spacemit::spine_mem_pool_shared_mem_alloc(size, alignment);
if (result == nullptr) {
GGML_LOG_ERROR("CPU_RISCV64_SPACEMIT: %s: failed to allocate shared memory size %zu alignment %zu\n", __func__,
size, alignment);
}
return result;
}
void ggml_backend_cpu_riscv64_spacemit_free_shared(void * ptr) {
ggml::cpu::riscv64_spacemit::spine_mem_pool_shared_mem_free(ptr);
}
}

32
ggml/src/ggml-cpu/spacemit/spine_mem_pool.h Normal file
View file

@ -0,0 +1,32 @@
#pragma once
#include <cstddef>
#include <cstdint>
namespace ggml::cpu::riscv64_spacemit {
enum class spine_mem_pool_backend : uint8_t {
none,
posix_memalign,
transparent_hugepage,
hugetlb_1g,
};
struct spine_mem_pool_tcm_info {
bool available{ false };
size_t blk_size{ 0 };
size_t blk_num{ 0 };
bool is_fake_tcm{ false };
};
bool spine_mem_pool_tcm_init(spine_mem_pool_tcm_info * info) noexcept;
void * spine_mem_pool_tcm_mem_get(int cpu_id) noexcept;
void * spine_mem_pool_tcm_mem_wait(int cpu_id) noexcept;
int spine_mem_pool_tcm_mem_release(int cpu_id) noexcept;
void * spine_mem_pool_alloc(size_t size, size_t alignment) noexcept;
void * spine_mem_pool_shared_mem_alloc(size_t size, size_t alignment) noexcept;
void spine_mem_pool_free(void * base) noexcept;
void spine_mem_pool_shared_mem_free(void * base) noexcept;
} // namespace ggml::cpu::riscv64_spacemit

409
ggml/src/ggml-cpu/spacemit/spine_tcm.h Normal file
View file

@ -0,0 +1,409 @@
#ifndef SPINE_TCM_PUBLIC_H_
#define SPINE_TCM_PUBLIC_H_
/*
* spine_tcm public API
*
* Usage:
* 1. Direct link mode
* Define SPINE_TCM_DIRECT_LINK and link against libspine_tcm.so.
*
* if (spine_tcm_is_available()) {
* void *buffer = spine_tcm_mem_get(0);
* spine_tcm_mem_free(0);
* }
*
* 2. Header-only loader mode
* Include this header without linking libspine_tcm.so. The loader first
* tries to reuse a process-global spine_tcm instance and falls back to
* dlopen("libspine_tcm.so") when needed.
*
* spine_tcm_open_handle(NULL); // optional pre-bind
* if (spine_tcm_is_available()) {
* void *buffer = spine_tcm_mem_get(0);
* spine_tcm_mem_free(0);
* }
*/
#include <stddef.h>
#include <stdint.h>
#include <string.h>
#if !defined(SPINE_TCM_BUILD_SHARED) && !defined(SPINE_TCM_DIRECT_LINK)
# include <dlfcn.h>
#endif
#ifdef __cplusplus
extern "C" {
#endif
#if defined(_WIN32)
# if defined(SPINE_TCM_BUILD_SHARED)
# define SPINE_TCM_API __declspec(dllexport)
# else
# define SPINE_TCM_API __declspec(dllimport)
# endif
#else
# define SPINE_TCM_API __attribute__((visibility("default")))
#endif
typedef struct spine_tcm_mem_info {
size_t blk_size;
size_t blk_num;
int is_fake_tcm;
} spine_tcm_mem_info_t;
typedef struct spine_tcm_block_info {
int id;
void * va;
size_t size;
uint64_t phys_addr;
uint64_t cpu_affinity_mask;
int owner_tid;
int is_acquired;
} spine_tcm_block_info_t;
/* Shared-library runtime ABI exported by libspine_tcm.so. */
SPINE_TCM_API const char * spine_tcm_runtime_version(void);
SPINE_TCM_API int spine_tcm_runtime_is_available(void);
SPINE_TCM_API int spine_tcm_runtime_layout_info(spine_tcm_mem_info_t * info);
SPINE_TCM_API int spine_tcm_runtime_mem_info(int id, spine_tcm_block_info_t * info);
SPINE_TCM_API void * spine_tcm_runtime_mem_get(int id);
SPINE_TCM_API int spine_tcm_runtime_mem_free(int id);
SPINE_TCM_API void * spine_tcm_runtime_mem_try_wait(int id, size_t timeout_us);
SPINE_TCM_API int spine_tcm_runtime_mem_release(int id);
SPINE_TCM_API int spine_tcm_runtime_mem_force_release(int id);
SPINE_TCM_API int spine_tcm_runtime_mem_query(int id);
#if defined(SPINE_TCM_DIRECT_LINK)
/* Optional no-op in direct-link mode. */
static inline int spine_tcm_open_handle(const char * so_path) {
(void) so_path;
return 0;
}
static inline const char * spine_tcm_version(void) {
return spine_tcm_runtime_version();
}
/* Returns 1 when the runtime driver is available, otherwise 0. */
static inline int spine_tcm_is_available(void) {
return spine_tcm_runtime_is_available();
}
/* Returns runtime memory geometry and whether the current backend is fake TCM. */
static inline int spine_tcm_mem_info(spine_tcm_mem_info_t * info) {
return spine_tcm_runtime_layout_info(info);
}
/* Returns per-block runtime metadata for the given TCM id. */
static inline int spine_tcm_block_info(int id, spine_tcm_block_info_t * info) {
return spine_tcm_runtime_mem_info(id, info);
}
/* Returns a cached buffer for the given TCM id, or NULL on failure. */
static inline void * spine_tcm_mem_get(int id) {
return spine_tcm_runtime_mem_get(id);
}
/* Releases one reference acquired by spine_tcm_mem_get(id). */
static inline int spine_tcm_mem_free(int id) {
return spine_tcm_runtime_mem_free(id);
}
/* Waits for a TCM block handoff and returns the driver-owned buffer when available. */
static inline void * spine_tcm_mem_try_wait(int id, size_t over_time) {
return spine_tcm_runtime_mem_try_wait(id, over_time);
}
/* Releases a buffer acquired by spine_tcm_mem_try_wait(id, over_time). */
static inline int spine_tcm_mem_release(int id) {
return spine_tcm_runtime_mem_release(id);
}
/* Forces a release for the given TCM id when the backend supports it. */
static inline int spine_tcm_mem_force_release(int id) {
return spine_tcm_runtime_mem_force_release(id);
}
/* Returns whether the given TCM id is currently acquired. */
static inline int spine_tcm_mem_query(int id) {
return spine_tcm_runtime_mem_query(id);
}
#elif !defined(SPINE_TCM_BUILD_SHARED)
typedef struct spine_tcm_handle {
void * module_handle;
int use_global_scope;
int owns_module_handle;
const char * (*runtime_version)(void);
int (*runtime_is_available)(void);
int (*runtime_layout_info)(spine_tcm_mem_info_t * info);
int (*runtime_mem_info)(int id, spine_tcm_block_info_t * info);
void * (*runtime_mem_get)(int id);
int (*runtime_mem_free)(int id);
void * (*runtime_mem_try_wait)(int id, size_t over_time);
int (*runtime_mem_release)(int id);
int (*runtime_mem_force_release)(int id);
int (*runtime_mem_query)(int id);
} spine_tcm_handle_t;
static inline spine_tcm_handle_t * spine_tcm_default_handle(void) {
static spine_tcm_handle_t handle = { 0 };
return &handle;
}
static inline void spine_tcm_handle_reset(spine_tcm_handle_t * handle) {
if (handle != NULL) {
memset(handle, 0, sizeof(*handle));
}
}
static inline int spine_tcm_handle_bind(spine_tcm_handle_t * handle) {
void * symbol_scope = handle->use_global_scope ? RTLD_DEFAULT : handle->module_handle;
handle->runtime_version = (const char * (*) (void) ) dlsym(symbol_scope, "spine_tcm_runtime_version");
handle->runtime_is_available = (int (*)(void)) dlsym(symbol_scope, "spine_tcm_runtime_is_available");
handle->runtime_layout_info =
(int (*)(spine_tcm_mem_info_t *)) dlsym(symbol_scope, "spine_tcm_runtime_layout_info");
handle->runtime_mem_info =
(int (*)(int, spine_tcm_block_info_t *)) dlsym(symbol_scope, "spine_tcm_runtime_mem_info");
handle->runtime_mem_get = (void * (*) (int) ) dlsym(symbol_scope, "spine_tcm_runtime_mem_get");
handle->runtime_mem_free = (int (*)(int)) dlsym(symbol_scope, "spine_tcm_runtime_mem_free");
handle->runtime_mem_try_wait = (void * (*) (int, size_t)) dlsym(symbol_scope, "spine_tcm_runtime_mem_try_wait");
handle->runtime_mem_release = (int (*)(int)) dlsym(symbol_scope, "spine_tcm_runtime_mem_release");
handle->runtime_mem_force_release = (int (*)(int)) dlsym(symbol_scope, "spine_tcm_runtime_mem_force_release");
handle->runtime_mem_query = (int (*)(int)) dlsym(symbol_scope, "spine_tcm_runtime_mem_query");
return handle->runtime_version != NULL && handle->runtime_is_available != NULL &&
handle->runtime_layout_info != NULL && handle->runtime_mem_info != NULL &&
handle->runtime_mem_get != NULL && handle->runtime_mem_free != NULL &&
handle->runtime_mem_try_wait != NULL && handle->runtime_mem_release != NULL &&
handle->runtime_mem_force_release != NULL && handle->runtime_mem_query != NULL ?
0 :
-1;
}
/*
* Try to bind against an already-loaded process-global spine_tcm instance.
* The shared library exports spine_tcm_runtime_marker only for this probe.
*/
static inline int spine_tcm_try_bind_global(spine_tcm_handle_t * handle) {
if (dlsym(RTLD_DEFAULT, "spine_tcm_runtime_marker") == NULL) {
return -1;
}
handle->use_global_scope = 1;
return spine_tcm_handle_bind(handle);
}
/*
* Optional pre-bind entry point.
*
* Behavior:
* - Reuses an already-loaded global spine_tcm instance when available.
* - Otherwise loads the shared library from so_path or the default soname.
* - Repeated calls are safe and return 0 after the first successful bind.
*/
static inline int spine_tcm_open_handle(const char * so_path) {
spine_tcm_handle_t * resolved = spine_tcm_default_handle();
const char * library = (so_path != NULL && so_path[0] != '\0') ? so_path : "libspine_tcm.so";
if (resolved->module_handle != NULL || resolved->use_global_scope) {
return 0;
}
if (spine_tcm_try_bind_global(resolved) == 0) {
return 0;
}
spine_tcm_handle_reset(resolved);
resolved->module_handle = dlopen(library, RTLD_LAZY | RTLD_GLOBAL);
resolved->owns_module_handle = resolved->module_handle != NULL ? 1 : 0;
if (resolved->module_handle == NULL) {
spine_tcm_handle_reset(resolved);
return -1;
}
if (spine_tcm_handle_bind(resolved) != 0) {
if (resolved->owns_module_handle) {
dlclose(resolved->module_handle);
}
spine_tcm_handle_reset(resolved);
return -1;
}
return 0;
}
/* Returns 1 when the runtime driver is available, otherwise 0. */
static inline int spine_tcm_is_available(void) {
spine_tcm_handle_t * resolved = spine_tcm_default_handle();
if (resolved->module_handle == NULL && !resolved->use_global_scope) {
(void) spine_tcm_open_handle(NULL);
}
if ((resolved->module_handle == NULL && !resolved->use_global_scope) || resolved->runtime_is_available == NULL) {
return 0;
}
return resolved->runtime_is_available();
}
/* Returns runtime memory geometry and whether the current backend is fake TCM. */
static inline int spine_tcm_mem_info(spine_tcm_mem_info_t * info) {
spine_tcm_handle_t * resolved = spine_tcm_default_handle();
if (resolved->module_handle == NULL && !resolved->use_global_scope) {
(void) spine_tcm_open_handle(NULL);
}
if ((resolved->module_handle == NULL && !resolved->use_global_scope) || resolved->runtime_layout_info == NULL) {
return -1;
}
return resolved->runtime_layout_info(info);
}
static inline const char * spine_tcm_version(void) {
spine_tcm_handle_t * resolved = spine_tcm_default_handle();
if (resolved->module_handle == NULL && !resolved->use_global_scope) {
(void) spine_tcm_open_handle(NULL);
}
if ((resolved->module_handle == NULL && !resolved->use_global_scope) || resolved->runtime_version == NULL) {
return "unknown";
}
return resolved->runtime_version();
}
/* Returns per-block runtime metadata for the given TCM id. */
static inline int spine_tcm_block_info(int id, spine_tcm_block_info_t * info) {
spine_tcm_handle_t * resolved = spine_tcm_default_handle();
if (resolved->module_handle == NULL && !resolved->use_global_scope) {
(void) spine_tcm_open_handle(NULL);
}
if ((resolved->module_handle == NULL && !resolved->use_global_scope) || resolved->runtime_mem_info == NULL) {
return -1;
}
return resolved->runtime_mem_info(id, info);
}
/* Returns a cached buffer for the given TCM id, or NULL on failure. */
static inline void * spine_tcm_mem_get(int id) {
spine_tcm_handle_t * resolved = spine_tcm_default_handle();
if (resolved->module_handle == NULL && !resolved->use_global_scope) {
(void) spine_tcm_open_handle(NULL);
}
if (resolved->module_handle == NULL && !resolved->use_global_scope) {
return NULL;
}
if (resolved->runtime_mem_get == NULL) {
return NULL;
}
return resolved->runtime_mem_get(id);
}
/* Releases one reference acquired by spine_tcm_mem_get(id). */
static inline int spine_tcm_mem_free(int id) {
spine_tcm_handle_t * resolved = spine_tcm_default_handle();
if (resolved->module_handle == NULL && !resolved->use_global_scope) {
(void) spine_tcm_open_handle(NULL);
}
if ((resolved->module_handle == NULL && !resolved->use_global_scope) || resolved->runtime_mem_free == NULL) {
return -1;
}
return resolved->runtime_mem_free(id);
}
/* Waits for a TCM block handoff and returns the driver-owned buffer when available. */
static inline void * spine_tcm_mem_try_wait(int id, size_t over_time) {
spine_tcm_handle_t * resolved = spine_tcm_default_handle();
if (resolved->module_handle == NULL && !resolved->use_global_scope) {
(void) spine_tcm_open_handle(NULL);
}
if (resolved->module_handle == NULL && !resolved->use_global_scope) {
return NULL;
}
if (resolved->runtime_mem_try_wait == NULL) {
return NULL;
}
return resolved->runtime_mem_try_wait(id, over_time);
}
/* Releases a buffer acquired by spine_tcm_mem_try_wait(id, over_time). */
static inline int spine_tcm_mem_release(int id) {
spine_tcm_handle_t * resolved = spine_tcm_default_handle();
if (resolved->module_handle == NULL && !resolved->use_global_scope) {
(void) spine_tcm_open_handle(NULL);
}
if ((resolved->module_handle == NULL && !resolved->use_global_scope) || resolved->runtime_mem_release == NULL) {
return -1;
}
return resolved->runtime_mem_release(id);
}
/* Forces a release for the given TCM id when the backend supports it. */
static inline int spine_tcm_mem_force_release(int id) {
spine_tcm_handle_t * resolved = spine_tcm_default_handle();
if (resolved->module_handle == NULL && !resolved->use_global_scope) {
(void) spine_tcm_open_handle(NULL);
}
if ((resolved->module_handle == NULL && !resolved->use_global_scope) ||
resolved->runtime_mem_force_release == NULL) {
return -1;
}
return resolved->runtime_mem_force_release(id);
}
/* Returns whether the given TCM id is currently acquired. */
static inline int spine_tcm_mem_query(int id) {
spine_tcm_handle_t * resolved = spine_tcm_default_handle();
if (resolved->module_handle == NULL && !resolved->use_global_scope) {
(void) spine_tcm_open_handle(NULL);
}
if ((resolved->module_handle == NULL && !resolved->use_global_scope) || resolved->runtime_mem_query == NULL) {
return -1;
}
return resolved->runtime_mem_query(id);
}
#else
static inline const char * spine_tcm_version(void) {
return spine_tcm_runtime_version();
}
#endif
#define SPINE_TCM_VERSION (spine_tcm_version())
#ifdef __cplusplus
}
#endif
#endif

11
ggml/src/ggml-cuda/allreduce.cu
View file

@ -184,13 +184,15 @@ static __global__ void ggml_cuda_ar_kernel(
#pragma unroll
for (int k = 0; k < ELEMS_PER_VEC; ++k) {
const T_wire d_low = ggml_cuda_cast<T_wire>(sendbuf[off + k]);
recvbuf[off + k] = ggml_cuda_cast<T_dst>(d_low) + ggml_cuda_cast<T_dst>(wire[k]);
recvbuf[off + k] = ggml_cuda_cast<T_dst>(
ggml_cuda_cast<float>(d_low) + ggml_cuda_cast<float>(wire[k]));
}
}
if (bid == 0 && tid < count - tail) {
const T_wire d_low = ggml_cuda_cast<T_wire>(sendbuf[tail + tid]);
recvbuf[tail + tid] =
ggml_cuda_cast<T_dst>(d_low) + ggml_cuda_cast<T_dst>(host_other[tail + tid]);
recvbuf[tail + tid] = ggml_cuda_cast<T_dst>(
ggml_cuda_cast<float>(d_low) +
ggml_cuda_cast<float>(host_other[tail + tid]));
}
}
}
@ -210,7 +212,8 @@ static __global__ void ggml_cuda_ar_add_kernel(
const int nt = gridDim.x * blockDim.x;
for (int i = tid; i < count; i += nt) {
const T_src d_low = ggml_cuda_cast<T_src>(dst[i]);
dst[i] = ggml_cuda_cast<T_dst>(d_low) + ggml_cuda_cast<T_dst>(src[i]);
dst[i] = ggml_cuda_cast<T_dst>(
ggml_cuda_cast<float>(d_low) + ggml_cuda_cast<float>(src[i]));
}
}

2
ggml/src/ggml-metal/ggml-metal-device.m
View file

@ -678,7 +678,7 @@ ggml_metal_device_t ggml_metal_device_init(int device) {
![[dev->mtl_device name] containsString:@"M6"] &&
![[dev->mtl_device name] containsString:@"A19"] &&
![[dev->mtl_device name] containsString:@"A20"]) {
GGML_LOG_WARN("%s: tensor API disabled for pre-M5 and pre-A19 devices\n", __func__);
GGML_LOG_INFO("%s: tensor API disabled for pre-M5 and pre-A19 devices\n", __func__);
dev->props.has_tensor = false;
}

24
ggml/src/ggml-vulkan/ggml-vulkan.cpp
View file

@ -3959,13 +3959,13 @@ static void ggml_vk_load_shaders(vk_device& device) {
ggml_vk_create_pipeline(device, device-> PIPELINE_NAME ->a_s, #NAMELC #F16ACC "_aligned_s", NAMELC ## _aligned ## F16ACC ## _len, NAMELC ## _aligned ## F16ACC ## _data, "main", PARAMCOUNT, sizeof(PUSHCONST), s_ ## WG_DENOMS, s_ ## WARPTILE, s_align, false, REQSUBGROUPSIZE > 0, REQSUBGROUPSIZE); \
#define CREATE_MMQ(TYPE, PIPELINE_NAME, NAMELC, WG_DENOMS, WARPTILE, PUSHCONST, PARAMCOUNT, ID, REQSUBGROUPSIZE) \
if (device->mul_mat ## ID ## _l[TYPE]) { \
if (device->mul_mat ## ID ## _l_int[TYPE]) { \
ggml_vk_create_pipeline(device, device-> PIPELINE_NAME .f32acc->l, #NAMELC "_l", NAMELC ## _len, NAMELC ## _data, "main", PARAMCOUNT, sizeof(PUSHCONST), l_ ## WG_DENOMS, l_ ## WARPTILE, 1, false, REQSUBGROUPSIZE > 0, REQSUBGROUPSIZE); \
} \
if (device->mul_mat ## ID ## _m[TYPE]) { \
if (device->mul_mat ## ID ## _m_int[TYPE]) { \
ggml_vk_create_pipeline(device, device-> PIPELINE_NAME .f32acc->m, #NAMELC "_m", NAMELC ## _len, NAMELC ## _data, "main", PARAMCOUNT, sizeof(PUSHCONST), m_ ## WG_DENOMS, m_ ## WARPTILE, 1, false, REQSUBGROUPSIZE > 0, REQSUBGROUPSIZE); \
} \
if (device->mul_mat ## ID ## _s[TYPE]) { \
if (device->mul_mat ## ID ## _s_int[TYPE]) { \
ggml_vk_create_pipeline(device, device-> PIPELINE_NAME .f32acc->s, #NAMELC "_s", NAMELC ## _len, NAMELC ## _data, "main", PARAMCOUNT, sizeof(PUSHCONST), s_ ## WG_DENOMS, s_ ## WARPTILE, 1, false, REQSUBGROUPSIZE > 0, REQSUBGROUPSIZE); \
} \
@ -4136,11 +4136,11 @@ static void ggml_vk_load_shaders(vk_device& device) {
ggml_vk_create_pipeline(device, device-> PIPELINE_NAME ->a_s, #NAMELC #F16ACC "_aligned_s", NAMELC ## _aligned ## F16ACC ## _fp32_len, NAMELC ## _aligned ## F16ACC ## _fp32_data, "main", PARAMCOUNT, sizeof(PUSHCONST), s_ ## WG_DENOMS, s_ ## WARPTILE, s_align, false, REQSUBGROUPSIZE > 0, REQSUBGROUPSIZE); \
#define CREATE_MMQ(TYPE, PIPELINE_NAME, NAMELC, WG_DENOMS, WARPTILE, PUSHCONST, PARAMCOUNT, ID) \
if (device->mul_mat ## ID ## _l[TYPE]) \
if (device->mul_mat ## ID ## _l_int[TYPE]) \
ggml_vk_create_pipeline(device, device-> PIPELINE_NAME ->l, #NAMELC "_l", NAMELC ## _fp32_len, NAMELC ## _fp32_data, "main", PARAMCOUNT, sizeof(PUSHCONST), l_ ## WG_DENOMS, l_ ## WARPTILE, 1); \
if (device->mul_mat ## ID ## _m[TYPE]) \
if (device->mul_mat ## ID ## _m_int[TYPE]) \
ggml_vk_create_pipeline(device, device-> PIPELINE_NAME ->m, #NAMELC "_m", NAMELC ## _fp32_len, NAMELC ## _fp32_data, "main", PARAMCOUNT, sizeof(PUSHCONST), m_ ## WG_DENOMS, m_ ## WARPTILE, 1); \
if (device->mul_mat ## ID ## _s[TYPE]) \
if (device->mul_mat ## ID ## _s_int[TYPE]) \
ggml_vk_create_pipeline(device, device-> PIPELINE_NAME ->s, #NAMELC "_s", NAMELC ## _fp32_len, NAMELC ## _fp32_data, "main", PARAMCOUNT, sizeof(PUSHCONST), s_ ## WG_DENOMS, s_ ## WARPTILE, 1); \
CREATE_MM(GGML_TYPE_F32, pipeline_matmul_f32, matmul_f32_f32, , wg_denoms, warptile, vk_mat_mat_push_constants, 3, , 0);
@ -5735,12 +5735,12 @@ static vk_device ggml_vk_get_device(size_t idx) {
break;
}
device->mul_mat_l_int[i] = true;
device->mul_mat_m_int[i] = true;
device->mul_mat_s_int[i] = true;
device->mul_mat_id_l_int[i] = true;
device->mul_mat_id_m_int[i] = true;
device->mul_mat_id_s_int[i] = true;
device->mul_mat_l_int[i] = device->mul_mat_l[i];
device->mul_mat_m_int[i] = device->mul_mat_m[i];
device->mul_mat_s_int[i] = device->mul_mat_s[i];
device->mul_mat_id_l_int[i] = device->mul_mat_id_l[i];
device->mul_mat_id_m_int[i] = device->mul_mat_id_m[i];
device->mul_mat_id_s_int[i] = device->mul_mat_id_s[i];
}

133
src/unicode.cpp
View file

@ -605,6 +605,136 @@ static std::vector<size_t> unicode_regex_split_custom_qwen2(const std::string &
return bpe_offsets;
}
// Qwen3.5 system regex: "(?i:'s|'t|'re|'ve|'m|'ll|'d)|[^\\r\\n\\p{L}\\p{N}]?[\\p{L}\\p{M}]+|\\p{N}| ?[^\\s\\p{L}\\p{M}\\p{N}]+[\\r\\n]*|\\s*[\\r\\n]+|\\s+(?!\\S)|\\s+"
// Compared to Qwen2, letter-runs also consume Unicode combining marks (\p{M}): [\p{L}\p{M}]+ instead of \p{L}+
static std::vector<size_t> unicode_regex_split_custom_qwen35(const std::string & text, const std::vector<size_t> & offsets) {
std::vector<size_t> bpe_offsets; // store the offset of each word
bpe_offsets.reserve(offsets.size()); // Reserve memory for the approximate size
const auto cpts = unicode_cpts_from_utf8(text);
size_t start = 0;
for (auto offset : offsets) {
const size_t offset_ini = start;
const size_t offset_end = start + offset;
assert(offset_end <= cpts.size());
start = offset_end;
static const uint32_t OUT_OF_RANGE = 0xFFFFFFFF;
auto _get_cpt = [&] (const size_t pos) -> uint32_t {
return (offset_ini <= pos && pos < offset_end) ? cpts[pos] : OUT_OF_RANGE;
};
auto _get_flags = [&] (const size_t pos) -> unicode_cpt_flags {
return (offset_ini <= pos && pos < offset_end) ? unicode_cpt_flags_from_cpt(cpts[pos]) : unicode_cpt_flags{};
};
size_t _prev_end = offset_ini;
auto _add_token = [&] (const size_t end) -> size_t {
assert(_prev_end <= end && end <= offset_end);
size_t len = end - _prev_end;
if (len > 0) {
bpe_offsets.push_back(len);
}
_prev_end = end;
return len;
};
for (size_t pos = offset_ini; pos < offset_end; /*pos++*/ ) {
const uint32_t cpt = _get_cpt(pos);
const auto flags = _get_flags(pos);
// regex: (?i:'s|'t|'re|'ve|'m|'ll|'d) // case insensitive
if (cpt == '\'' && pos+1 < offset_end) {
uint32_t cpt_next = unicode_tolower(_get_cpt(pos+1));
if (cpt_next == 's' || cpt_next == 't' || cpt_next == 'm' || cpt_next == 'd') {
pos += _add_token(pos+2);
continue;
}
if (pos+2 < offset_end) {
uint32_t cpt_next_next = unicode_tolower(_get_cpt(pos+2));
if ((cpt_next == 'r' && cpt_next_next == 'e') ||
(cpt_next == 'v' && cpt_next_next == 'e') ||
(cpt_next == 'l' && cpt_next_next == 'l')) {
pos += _add_token(pos+3);
continue;
}
}
}
// regex: [^\r\n\p{L}\p{N}]?[\p{L}\p{M}]+
if (!(cpt == '\r' || cpt == '\n' || flags.is_number)) {
if (flags.is_letter || flags.is_accent_mark || _get_flags(pos + 1).is_accent_mark || _get_flags(pos+1).is_letter) {
pos++;
while (_get_flags(pos).is_letter || _get_flags(pos).is_accent_mark) {
pos++;
}
_add_token(pos);
continue;
}
}
// regex: \p{N}
if (flags.is_number) {
pos++;
_add_token(pos);
continue;
}
// regex: <space>?[^\s\p{L}\p{M}\p{N}]+[\r\n]*
auto flags2 = (cpt == ' ' ? _get_flags(pos+1) : flags);
if (!(flags2.is_whitespace | flags2.is_letter | flags2.is_accent_mark | flags2.is_number) && flags.as_uint()) {
pos += (cpt == ' ');
while (!(flags2.is_whitespace | flags2.is_letter | flags2.is_accent_mark | flags2.is_number) && flags2.as_uint()) {
flags2 = _get_flags(++pos);
}
uint32_t cpt2 = _get_cpt(pos);
while (cpt2 == '\r' || cpt2 == '\n') {
cpt2 = _get_cpt(++pos);
}
_add_token(pos);
continue;
}
size_t num_whitespaces = 0;
size_t last_end_r_or_n = 0;
while (_get_flags(pos+num_whitespaces).is_whitespace) {
uint32_t cpt2 = _get_cpt(pos+num_whitespaces);
if (cpt2 == '\r' || cpt2 == '\n') {
last_end_r_or_n = pos + num_whitespaces + 1;
}
num_whitespaces++;
}
// regex: \s*[\r\n]+
if (last_end_r_or_n > 0) {
pos = last_end_r_or_n;
_add_token(pos);
continue;
}
// regex: \s+(?!\S)
if (num_whitespaces > 1 && _get_cpt(pos+num_whitespaces) != OUT_OF_RANGE) {
pos += num_whitespaces - 1;
_add_token(pos);
continue;
}
// regex: \s+
if (num_whitespaces > 0) {
pos += num_whitespaces;
_add_token(pos);
continue;
}
// no matches
_add_token(++pos);
}
}
return bpe_offsets;
}
template <typename CharT>
static std::vector<size_t> unicode_regex_split_stl(const std::basic_string<CharT> & text, const std::basic_string<CharT> & regex, const std::vector<size_t> & offsets) {
using BidirIt = typename std::basic_string<CharT>::const_iterator;
@ -929,6 +1059,9 @@ static std::vector<size_t> unicode_regex_split_custom(const std::string & text,
} else if (
regex_expr == "(?:'[sS]|'[tT]|'[rR][eE]|'[vV][eE]|'[mM]|'[lL][lL]|'[dD])|[^\\r\\n\\p{L}\\p{N}]?\\p{L}+|\\p{N}| ?[^\\s\\p{L}\\p{N}]+[\\r\\n]*|\\s*[\\r\\n]+|\\s+(?!\\S)|\\s+") {
bpe_offsets = unicode_regex_split_custom_qwen2(text, offsets);
} else if (
regex_expr == "(?:'[sS]|'[tT]|'[rR][eE]|'[vV][eE]|'[mM]|'[lL][lL]|'[dD])|[^\\r\\n\\p{L}\\p{N}]?[\\p{L}\\p{M}]+|\\p{N}| ?[^\\s\\p{L}\\p{M}\\p{N}]+[\\r\\n]*|\\s*[\\r\\n]+|\\s+(?!\\S)|\\s+") {
bpe_offsets = unicode_regex_split_custom_qwen35(text, offsets);
} else if (regex_expr == "\\p{Han}+") {
// K2's first pattern - handle all K2 patterns together
bpe_offsets = unicode_regex_split_custom_kimi_k2(text, offsets);

2
tools/mtmd/clip-impl.h
View file

@ -474,10 +474,10 @@ static void clip_log_internal(enum ggml_log_level level, const char * format, ..
}
#ifndef LOG_INF
#define LOG_DBG(...) clip_log_internal(GGML_LOG_LEVEL_DEBUG, __VA_ARGS__)
#define LOG_INF(...) clip_log_internal(GGML_LOG_LEVEL_INFO, __VA_ARGS__)
#define LOG_WRN(...) clip_log_internal(GGML_LOG_LEVEL_WARN, __VA_ARGS__)
#define LOG_ERR(...) clip_log_internal(GGML_LOG_LEVEL_ERROR, __VA_ARGS__)
#define LOG_DBG(...) clip_log_internal(GGML_LOG_LEVEL_DEBUG, __VA_ARGS__)
#define LOG_CNT(...) clip_log_internal(GGML_LOG_LEVEL_CONT, __VA_ARGS__)
#endif

12
tools/mtmd/clip.cpp
View file

@ -1054,7 +1054,7 @@ struct clip_model_loader {
bool has_audio = false;
// TODO @ngxson : we should not pass clip_ctx here, it should be clip_model
clip_model_loader(const char * fname) : fname(fname) {
clip_model_loader(const char * fname, bool skip_tensors = false) : fname(fname) {
struct ggml_context * meta = nullptr;
struct gguf_init_params params = {
@ -1100,7 +1100,7 @@ struct clip_model_loader {
}
// tensors
{
if (!skip_tensors) {
for (int i = 0; i < n_tensors; ++i) {
const char * name = gguf_get_tensor_name(ctx_gguf.get(), i);
const size_t offset = gguf_get_tensor_offset(ctx_gguf.get(), i);
@ -3020,6 +3020,14 @@ struct clip_init_result clip_init(const char * fname, struct clip_context_params
return {ctx_vision, ctx_audio};
}
struct clip_cap clip_get_cap(const char * fname) {
clip_cap res;
clip_model_loader loader(fname, /* skip_tensors= */ true);
res.has_vision = loader.has_vision;
res.has_audio = loader.has_audio;
return res;
}
struct clip_image_size * clip_image_size_init() {
struct clip_image_size * load_image_size = new struct clip_image_size();
load_image_size->width = 448;

8
tools/mtmd/clip.h
View file

@ -126,4 +126,12 @@ bool clip_has_vision_encoder(const struct clip_ctx * ctx);
bool clip_has_audio_encoder(const struct clip_ctx * ctx);
bool clip_has_whisper_encoder(const struct clip_ctx * ctx);
bool clip_model_quantize(const char * fname_inp, const char * fname_out, const int itype) ;
struct clip_cap {
bool has_vision;
bool has_audio;
};
struct clip_cap clip_get_cap(const char * fname);

13
tools/mtmd/mtmd.cpp
View file

@ -1423,6 +1423,19 @@ void mtmd_log_set(ggml_log_callback log_callback, void * user_data) {
g_logger_state.log_callback_user_data = user_data;
}
struct mtmd_caps mtmd_get_cap_from_file(const char * fname) {
try {
auto tmp = clip_get_cap(fname);
mtmd_caps cap;
cap.inp_audio = tmp.has_audio;
cap.inp_vision = tmp.has_vision;
return cap;
} catch (const std::exception & e) {
LOG_ERR("%s: failed to get capabilities from file '%s': %s\n", __func__, fname, e.what());
return mtmd_caps{ false, false };
}
}
//
// Debugging API (NOT intended for public use)
//

8
tools/mtmd/mtmd.h
View file

@ -244,6 +244,14 @@ MTMD_API float * mtmd_get_output_embd(mtmd_context * ctx);
// If this is not called, or NULL is supplied, everything is output on stderr.
MTMD_API void mtmd_log_set(ggml_log_callback log_callback, void * user_data);
// EXPERIMENTAL API to get mmproj's capabilities without initializing the full context
// This is only intended to be used by llama-server, breaking changes is expected
struct mtmd_caps {
bool inp_vision;
bool inp_audio;
};
MTMD_API struct mtmd_caps mtmd_get_cap_from_file(const char * mmproj_fname);
/////////////////////////////////////////
// test function, to be used in test-mtmd-c-api.c

2
tools/server/public/bundle.css
View file

File diff suppressed because one or more lines are too long

6420
tools/server/public/bundle.js
View file

File diff suppressed because it is too large Load diff

56
tools/server/server-common.cpp
View file

@ -1040,6 +1040,10 @@ json oaicompat_chat_params_parse(
inputs.use_jinja = opt.use_jinja;
inputs.parallel_tool_calls = json_value(body, "parallel_tool_calls", caps["supports_parallel_tool_calls"]);
inputs.add_generation_prompt = json_value(body, "add_generation_prompt", true);
const bool continue_final_message = json_value(body, "continue_final_message", false);
if (continue_final_message && inputs.add_generation_prompt) {
throw std::invalid_argument("Cannot set both add_generation_prompt and continue_final_message to true.");
}
inputs.reasoning_format = opt.reasoning_format;
if (body.contains("reasoning_format")) {
inputs.reasoning_format = common_reasoning_format_from_name(body.at("reasoning_format").get<std::string>());
@ -1071,7 +1075,10 @@ json oaicompat_chat_params_parse(
// if the assistant message appears at the end of list, we do not add end-of-turn token
// for ex. this can be useful to modify the reasoning process in reasoning models
bool prefill_assistant_message = !inputs.messages.empty() && inputs.messages.back().role == "assistant" && opt.prefill_assistant;
// continue_final_message is the explicit opt in alias from the vLLM/transformers API,
// equivalent to the prefill_assistant heuristic
bool prefill_assistant_message = !inputs.messages.empty() && inputs.messages.back().role == "assistant"
&& (continue_final_message || opt.prefill_assistant);
common_chat_msg last_message;
if (prefill_assistant_message) {
last_message = inputs.messages.back();
@ -1082,11 +1089,12 @@ json oaicompat_chat_params_parse(
throw std::invalid_argument("Cannot have 2 or more assistant messages at the end of the list.");
}
/* TODO: test this properly */
inputs.reasoning_format = COMMON_REASONING_FORMAT_NONE;
if ( inputs.enable_thinking ) {
throw std::invalid_argument("Assistant response prefill is incompatible with enable_thinking.");
// reject reasoning prefill on channel based templates that do not expose explicit thinking tags
if (!last_message.reasoning_content.empty() && inputs.enable_thinking) {
auto probe_params = common_chat_templates_apply(opt.tmpls.get(), inputs);
if (probe_params.supports_thinking && probe_params.thinking_end_tag.empty()) {
throw std::invalid_argument("Assistant prefill with reasoning_content is not supported yet for this template.");
}
}
inputs.add_generation_prompt = true;
@ -1098,6 +1106,42 @@ json oaicompat_chat_params_parse(
/* Append assistant prefilled message */
if (prefill_assistant_message) {
const bool thinking_active = chat_params.supports_thinking && !chat_params.thinking_end_tag.empty();
const bool has_reasoning = !last_message.reasoning_content.empty();
const bool has_content = !last_message.content.empty() || !last_message.content_parts.empty();
const bool mid_reasoning = has_reasoning && !has_content;
// some templates inject thinking_start in generation_prompt, others let the model emit it
const bool gp_has_think = thinking_active
&& chat_params.generation_prompt.find(chat_params.thinking_start_tag) != std::string::npos;
// open the thinking block when reasoning is present and the template did not inject it
if (has_reasoning) {
if (thinking_active && !gp_has_think) {
chat_params.prompt += chat_params.thinking_start_tag;
}
chat_params.prompt += last_message.reasoning_content;
}
if (thinking_active) {
if (mid_reasoning) {
// model continues inside the thinking block, keep generation_prompt open on think
if (!gp_has_think) {
chat_params.generation_prompt += chat_params.thinking_start_tag;
}
} else {
// close thinking block when reasoning is followed by content, or when the template forced it open
if (has_reasoning || gp_has_think) {
chat_params.prompt += chat_params.thinking_end_tag;
}
// strip thinking_start from generation_prompt so the parser routes model output as content
auto pos = chat_params.generation_prompt.rfind(chat_params.thinking_start_tag);
if (pos != std::string::npos) {
chat_params.generation_prompt = chat_params.generation_prompt.substr(0, pos);
}
}
}
if (!last_message.content_parts.empty()) {
for (auto & p : last_message.content_parts) {
chat_params.prompt += p.text;

10
tools/server/server-common.h
View file

@ -15,17 +15,19 @@
using json = nlohmann::ordered_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__)
#define SLT_INF(slot, fmt, ...) LOG_INF("slot %12.*s: id %2d | task %d | " fmt, 12, __func__, (slot).id, ((slot).task ? (slot).task->id : -1), __VA_ARGS__)
#define SLT_CNT(slot, fmt, ...) LOG_CNT("" fmt, __VA_ARGS__)
#define SLT_WRN(slot, fmt, ...) LOG_WRN("slot %12.*s: id %2d | task %d | " fmt, 12, __func__, (slot).id, ((slot).task ? (slot).task->id : -1), __VA_ARGS__)
#define SLT_ERR(slot, fmt, ...) LOG_ERR("slot %12.*s: id %2d | task %d | " fmt, 12, __func__, (slot).id, ((slot).task ? (slot).task->id : -1), __VA_ARGS__)
#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_CNT(slot, fmt, ...) LOG_CNT("" fmt, __VA_ARGS__)
#define SRV_DBG(fmt, ...) LOG_DBG("srv %12.*s: " fmt, 12, __func__, __VA_ARGS__)
#define SRV_TRC(fmt, ...) LOG_TRC("srv %12.*s: " fmt, 12, __func__, __VA_ARGS__)
#define SRV_INF(fmt, ...) LOG_INF("srv %12.*s: " fmt, 12, __func__, __VA_ARGS__)
#define SRV_CNT(fmt, ...) LOG_CNT("" fmt, __VA_ARGS__)
#define SRV_WRN(fmt, ...) LOG_WRN("srv %12.*s: " fmt, 12, __func__, __VA_ARGS__)
#define SRV_ERR(fmt, ...) LOG_ERR("srv %12.*s: " fmt, 12, __func__, __VA_ARGS__)
#define SRV_DBG(fmt, ...) LOG_DBG("srv %12.*s: " fmt, 12, __func__, __VA_ARGS__)
#define SRV_CNT(fmt, ...) LOG_CNT("" fmt, __VA_ARGS__)
using raw_buffer = std::vector<uint8_t>;

87
tools/server/server-context.cpp
View file

@ -166,6 +166,7 @@ struct server_slot {
// stats
size_t n_sent_text = 0; // number of sent text character
int64_t t_print_last = 0;
int64_t t_start_process_prompt;
int64_t t_start_generation;
@ -233,7 +234,7 @@ struct server_slot {
}
}
SLT_INF(*this, "init sampler, took %0.2f ms, tokens: text = %d, total = %d\n",
SLT_TRC(*this, "init sampler, took %0.2f ms, tokens: text = %d, total = %d\n",
(ggml_time_us() - t_start) / 1000.0, n_text, (int) prompt.tokens.size());
}
@ -417,6 +418,36 @@ struct server_slot {
return stop_pos;
}
void print_timings_tg() {
if (n_decoded < 100) {
return;
}
const int64_t t_now = ggml_time_us();
if (t_now - t_print_last < 3*1000*1000) {
return;
}
t_print_last = t_now;
const double n_gen_second = 1e3 / t_token_generation * n_decoded;
SLT_INF(*this, "n_decoded = %6d, tg = %6.2f t/s\n", n_decoded, n_gen_second);
}
void print_timings_pp() const {
const double n_prompt_second = 1e3 / t_prompt_processing * n_prompt_tokens_processed;
const double f_progress = (float) prompt.n_tokens() / task->n_tokens();
if (t_prompt_processing < 3000.0) {
return;
}
SLT_INF(*this, "prompt processing, n_tokens = %6d, progress = %.2f, t = %6.2f s / %.2f tokens per second\n",
n_prompt_tokens_processed, f_progress, t_prompt_processing / 1e3, n_prompt_second);
}
void print_timings() const {
const double t_prompt = t_prompt_processing / n_prompt_tokens_processed;
const double n_prompt_second = 1e3 / t_prompt_processing * n_prompt_tokens_processed;
@ -588,6 +619,10 @@ public:
// note: chat_params must not be refreshed upon existing sleeping state
server_chat_params chat_params;
server_context_impl() {
mtmd_helper_log_set(common_log_default_callback, nullptr);
}
~server_context_impl() {
if (!sleeping) {
// destroy() is already called when entering sleeping state
@ -749,10 +784,6 @@ private:
std::string & mmproj_path = params_base.mmproj.path;
if (!mmproj_path.empty()) {
if (!is_resume) {
mtmd_helper_log_set(common_log_default_callback, nullptr);
}
mtmd_context_params mparams = mtmd_context_params_default();
mparams.use_gpu = params_base.mmproj_use_gpu;
@ -896,17 +927,17 @@ private:
if (params_base.cache_ram_mib != 0) {
if (params_base.cache_ram_mib < 0) {
SRV_WRN("prompt cache is enabled, size limit: %s\n", "no limit");
SRV_INF("prompt cache is enabled, size limit: %s\n", "no limit");
} else {
SRV_WRN("prompt cache is enabled, size limit: %d MiB\n", params_base.cache_ram_mib);
SRV_INF("prompt cache is enabled, size limit: %d MiB\n", params_base.cache_ram_mib);
}
SRV_WRN("%s", "use `--cache-ram 0` to disable the prompt cache\n");
SRV_INF("%s", "use `--cache-ram 0` to disable the prompt cache\n");
prompt_cache = std::make_unique<server_prompt_cache>(params_base.cache_ram_mib, n_ctx);
} else {
SRV_WRN("%s", "prompt cache is disabled - use `--cache-ram N` to enable it\n");
SRV_INF("%s", "prompt cache is disabled - use `--cache-ram N` to enable it\n");
}
SRV_WRN("%s", "for more info see https://github.com/ggml-org/llama.cpp/pull/16391\n");
SRV_INF("%s", "for more info see https://github.com/ggml-org/llama.cpp/pull/16391\n");
if (!params_base.model_alias.empty()) {
// backward compat: use first alias as model name
@ -954,13 +985,13 @@ private:
if (params_base.cache_idle_slots) {
if (!params_base.kv_unified) {
SRV_WRN("%s: --cache-idle-slots requires --kv-unified, disabling\n", __func__);
SRV_WRN("%s", "--cache-idle-slots requires --kv-unified, disabling\n");
params_base.cache_idle_slots = false;
} else if (params_base.cache_ram_mib == 0) {
SRV_WRN("%s: --cache-idle-slots requires --cache-ram, disabling\n", __func__);
SRV_WRN("%s", "--cache-idle-slots requires --cache-ram, disabling\n");
params_base.cache_idle_slots = false;
} else {
SRV_INF("%s: idle slots will be saved to prompt cache and cleared upon starting a new task\n", __func__);
SRV_INF("%s", "idle slots will be saved to prompt cache and cleared upon starting a new task\n");
SRV_DBG("%s", "__TEST_TAG_CACHE_IDLE_SLOTS_ENABLED__\n");
}
}
@ -1112,7 +1143,7 @@ private:
update_cache = update_cache && task.type == SERVER_TASK_TYPE_COMPLETION;
if (update_cache) {
SRV_WRN("%s", "updating prompt cache\n");
SRV_INF("%s", "updating prompt cache\n");
const int64_t t_start = ggml_time_us();
@ -1127,7 +1158,7 @@ private:
prompt_cache->update();
SRV_WRN("prompt cache update took %.2f ms\n", (ggml_time_us() - t_start) / 1000.0);
SRV_INF("prompt cache update took %.2f ms\n", (ggml_time_us() - t_start) / 1000.0);
}
}
@ -1186,10 +1217,10 @@ private:
if (!are_lora_equal(task_loras, slot.lora)) {
// if lora has changed, check to see if the cache should be cleared
if (lora_should_clear_cache(slot.lora, task_loras)) {
SLT_INF(slot, "clearing cache for lora change. %zu loras -> %zu loras\n", slot.lora.size(), task.params.lora.size());
SLT_TRC(slot, "clearing cache for lora change. %zu loras -> %zu loras\n", slot.lora.size(), task.params.lora.size());
slot.prompt.tokens.clear();
} else {
SLT_INF(slot, "keeping cache for alora. %zu target loras\n", task_loras.size());
SLT_TRC(slot, "keeping cache for alora. %zu target loras\n", task_loras.size());
}
slot.lora = task_loras;
}
@ -1281,7 +1312,8 @@ private:
llama_set_sampler(ctx_tgt, slot.id, nullptr);
}
SLT_INF(slot, "sampler chain: %s\n", common_sampler_print(slot.smpl.get()).c_str());
SLT_TRC(slot, "sampler chain: %s\n", common_sampler_print(slot.smpl.get()).c_str());
SLT_TRC(slot, "sampler params: \n%s\n", task.params.sampling.print().c_str());
} else {
slot.smpl.reset();
}
@ -1800,7 +1832,7 @@ private:
cur.update_tgt(ctx_tgt, slot.id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
cur.update_dft(ctx_dft.get(), slot.id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
SLT_WRN(slot,
SLT_INF(slot,
"created context checkpoint %d of %d (pos_min = %d, pos_max = %d, n_tokens = %" PRId64 ", size = %.3f MiB)\n",
(int) slot.prompt.checkpoints.size(), params_base.n_ctx_checkpoints, cur.pos_min,
cur.pos_max, cur.n_tokens, (float) cur.size() / 1024 / 1024);
@ -2339,7 +2371,7 @@ private:
slot.state = SLOT_STATE_PROCESSING_PROMPT;
SLT_INF(slot, "new prompt, n_ctx_slot = %d, n_keep = %d, task.n_tokens = %d\n",
SLT_TRC(slot, "new prompt, n_ctx_slot = %d, n_keep = %d, task.n_tokens = %d\n",
slot.n_ctx, slot.task->params.n_keep, slot.task->n_tokens());
// print prompt tokens (for debugging)
@ -2454,7 +2486,7 @@ private:
}
if (n_match >= (size_t) n_cache_reuse) {
SLT_INF(slot, "reusing chunk with size %zu, shifting KV cache [%zu, %zu) -> [%zu, %zu)\n", n_match, head_c, head_c + n_match, head_p, head_p + n_match);
SLT_TRC(slot, "reusing chunk with size %zu, shifting KV cache [%zu, %zu) -> [%zu, %zu)\n", n_match, head_c, head_c + n_match, head_p, head_p + n_match);
//for (size_t i = head_p; i < head_p + n_match; i++) {
// SLT_DBG(slot, "cache token %3zu: %6d '%s'\n", i, prompt_tokens[i], common_token_to_piece(ctx_tgt, prompt_tokens[i]).c_str());
//}
@ -2620,10 +2652,14 @@ private:
}
}
const int64_t t_current = ggml_time_us();
slot.t_prompt_processing = (t_current - slot.t_start_process_prompt) / 1e3;
slot.print_timings_pp();
// truncate any tokens that are beyond n_past for this slot
const llama_pos p0 = slot.prompt.tokens.pos_next();
SLT_INF(slot, "n_tokens = %d, memory_seq_rm [%d, end)\n", slot.prompt.n_tokens(), p0);
SLT_TRC(slot, "cached n_tokens = %d, memory_seq_rm [%d, end)\n", slot.prompt.n_tokens(), p0);
if (!llama_memory_seq_rm(llama_get_memory(ctx_tgt), slot.id, p0, -1)) {
SLT_WRN(slot, "failed to truncate tokens with position >= %d - clearing the memory\n", p0);
@ -2764,7 +2800,6 @@ private:
slot.i_batch = batch.n_tokens - 1;
slot.init_sampler();
SLT_INF(slot, "prompt processing done, n_tokens = %d, batch.n_tokens = %d\n", slot.prompt.n_tokens(), batch.n_tokens);
} else {
if (slot.task->n_tokens() < slot.prompt.n_tokens() + n_ubatch) {
// near the end of the prompt
@ -2786,8 +2821,6 @@ private:
}
}
}
SLT_INF(slot, "prompt processing progress, n_tokens = %d, batch.n_tokens = %d, progress = %f\n", slot.prompt.n_tokens(), batch.n_tokens, (float) slot.prompt.n_tokens() / slot.task->n_tokens());
}
const auto pos_min = llama_memory_seq_pos_min(llama_get_memory(ctx_tgt), slot.id);
@ -3084,6 +3117,8 @@ private:
continue;
}
slot.print_timings_tg();
}
// speculative decoding - main model sample and accept
@ -3196,6 +3231,8 @@ private:
}
}
slot.print_timings_tg();
SLT_DBG(slot, "accepted %d/%d draft tokens, new n_tokens = %d\n", (int) ids.size() - 1, (int) n_draft, slot.prompt.n_tokens());
}
}

38
tools/server/server-http.cpp
View file

@ -47,7 +47,7 @@ static void log_server_request(const httplib::Request & req, const httplib::Resp
// reminder: this function is not covered by httplib's exception handler; if someone does more complicated stuff, think about wrapping it in try-catch
SRV_INF("done request: %s %s %s %d\n", req.method.c_str(), req.path.c_str(), req.remote_addr.c_str(), res.status);
SRV_TRC("done request: %s %s %s %d\n", req.method.c_str(), req.path.c_str(), req.remote_addr.c_str(), res.status);
SRV_DBG("request: %s\n", req.body.c_str());
SRV_DBG("response: %s\n", res.body.c_str());
@ -89,10 +89,10 @@ bool server_http_context::init(const common_params & params) {
hostname = params.hostname;
if (gcp.enabled) {
LOG_INF("%s: Google Cloud Platform compat: health route = %s, predict route = %s, port = %d\n", __func__, gcp.path_health.c_str(), gcp.path_predict.c_str(), gcp.port);
SRV_INF("Google Cloud Platform compat: health route = %s, predict route = %s, port = %d\n", gcp.path_health.c_str(), gcp.path_predict.c_str(), gcp.port);
if (port != gcp.port) {
LOG_WRN("%s: Google Cloud Platform compat: overriding server port %d with AIP_HTTP_PORT %d\n", __func__, port, gcp.port);
SRV_WRN("Google Cloud Platform compat: overriding server port %d with AIP_HTTP_PORT %d\n", port, gcp.port);
}
port = gcp.port;
@ -102,17 +102,17 @@ bool server_http_context::init(const common_params & params) {
#ifdef CPPHTTPLIB_OPENSSL_SUPPORT
if (params.ssl_file_key != "" && params.ssl_file_cert != "") {
LOG_INF("Running with SSL: key = %s, cert = %s\n", params.ssl_file_key.c_str(), params.ssl_file_cert.c_str());
SRV_INF("running with SSL: key = %s, cert = %s\n", params.ssl_file_key.c_str(), params.ssl_file_cert.c_str());
srv.reset(
new httplib::SSLServer(params.ssl_file_cert.c_str(), params.ssl_file_key.c_str())
);
} else {
LOG_INF("Running without SSL\n");
SRV_INF("%s", "running without SSL\n");
srv.reset(new httplib::Server());
}
#else
if (params.ssl_file_key != "" && params.ssl_file_cert != "") {
LOG_ERR("Server is built without SSL support\n");
SRV_ERR("%s", "the server is built without SSL support\n");
return false;
}
srv.reset(new httplib::Server());
@ -134,7 +134,7 @@ bool server_http_context::init(const common_params & params) {
res.status = 500;
res.set_content(message, "text/plain");
LOG_ERR("got exception: %s\n", message.c_str());
SRV_ERR("got exception: %s\n", message.c_str());
});
srv->set_error_handler([](const httplib::Request &, httplib::Response & res) {
@ -162,7 +162,7 @@ bool server_http_context::init(const common_params & params) {
#ifdef SO_REUSEPORT
httplib::set_socket_opt(sock, SOL_SOCKET, SO_REUSEPORT, 1);
#else
LOG_WRN("%s: SO_REUSEPORT is not supported\n", __func__);
SRV_WRN("%s", "SO_REUSEPORT is not supported\n");
#endif
}
});
@ -170,9 +170,9 @@ bool server_http_context::init(const common_params & params) {
if (params.api_keys.size() == 1) {
auto key = params.api_keys[0];
std::string substr = key.substr(std::max((int)(key.length() - 4), 0));
LOG_INF("%s: api_keys: ****%s\n", __func__, substr.c_str());
SRV_INF("api_keys: ****%s\n", substr.c_str());
} else if (params.api_keys.size() > 1) {
LOG_INF("%s: api_keys: %zu keys loaded\n", __func__, params.api_keys.size());
SRV_INF("api_keys: %zu keys loaded\n", params.api_keys.size());
}
//
@ -232,7 +232,7 @@ bool server_http_context::init(const common_params & params) {
"application/json; charset=utf-8"
);
LOG_WRN("Unauthorized: Invalid API Key\n");
SRV_WRN("%s", "unauthorized: Invalid API Key\n");
return false;
};
@ -292,7 +292,7 @@ bool server_http_context::init(const common_params & params) {
// +4 threads for monitoring, health and some threads reserved for MCP and other tasks in the future
n_threads_http = std::max(params.n_parallel + 4, (int32_t) std::thread::hardware_concurrency() - 1);
}
LOG_INF("%s: using %d threads for HTTP server\n", __func__, n_threads_http);
SRV_INF("using %d threads for HTTP server\n", n_threads_http);
srv->new_task_queue = [n_threads_http] {
// spawn n_threads_http fixed thread (always alive), while allow up to 1024 max possible additional threads
// when n_threads_http is used, server will create new "dynamic" threads that will be destroyed after processing each request
@ -306,14 +306,14 @@ bool server_http_context::init(const common_params & params) {
//
if (!params.webui) {
LOG_INF("Web UI is disabled\n");
SRV_INF("%s", "the WebUI is disabled\n");
} else {
// register static assets routes
if (!params.public_path.empty()) {
// Set the base directory for serving static files
bool is_found = srv->set_mount_point(params.api_prefix + "/", params.public_path);
if (!is_found) {
LOG_ERR("%s: static assets path not found: %s\n", __func__, params.public_path.c_str());
SRV_ERR("static assets path not found: %s\n", params.public_path.c_str());
return 1;
}
} else {
@ -348,13 +348,13 @@ bool server_http_context::start() {
bool is_sock = false;
if (string_ends_with(std::string(hostname), ".sock")) {
is_sock = true;
LOG_INF("%s: setting address family to AF_UNIX\n", __func__);
SRV_INF("%s", "setting address family to AF_UNIX\n");
srv->set_address_family(AF_UNIX);
// bind_to_port requires a second arg, any value other than 0 should
// simply get ignored
was_bound = srv->bind_to_port(hostname, 8080);
} else {
LOG_INF("%s: binding port with default address family\n", __func__);
SRV_INF("%s", "binding port with default address family\n");
// bind HTTP listen port
if (port == 0) {
int bound_port = srv->bind_to_any_port(hostname);
@ -368,7 +368,7 @@ bool server_http_context::start() {
}
if (!was_bound) {
LOG_ERR("%s: couldn't bind HTTP server socket, hostname: %s, port: %d\n", __func__, hostname.c_str(), port);
SRV_ERR("couldn't bind HTTP server socket, hostname: %s, port: %d\n", hostname.c_str(), port);
return false;
}
@ -580,7 +580,7 @@ void server_http_context::register_gcp_compat() {
}
if (handlers.count(gcp.path_predict)) {
LOG_ERR("%s: AIP_PREDICT_ROUTE=%s conflicts with an existing llama-server route\n", __func__, gcp.path_predict.c_str());
SRV_ERR("AIP_PREDICT_ROUTE=%s conflicts with an existing llama-server route\n", gcp.path_predict.c_str());
exit(1);
}
@ -651,7 +651,7 @@ void server_http_context::register_gcp_compat() {
payload.erase("@requestFormat");
if (payload.contains("stream")) {
LOG_WRN("%s: ignoring client-provided stream field in instance, streaming is not supported in predict route\n", __func__);
SRV_WRN("%s", "ignoring client-provided stream field in instance, streaming is not supported in predict route\n");
payload["stream"] = false;
}

58
tools/server/server-models.cpp
View file

@ -161,6 +161,30 @@ void server_model_meta::update_args(common_preset_context & ctx_preset, std::str
args = preset.to_args(bin_path);
}
void server_model_meta::update_caps() {
try {
common_params params;
preset.apply_to_params(params, {
"LLAMA_ARG_MODEL",
"LLAMA_ARG_MODEL_URL",
"LLAMA_ARG_MMPROJ",
"LLAMA_ARG_MMPROJ_URL",
"LLAMA_ARG_HF_REPO",
"LLAMA_ARG_HF_REPO_FILE",
});
params.offline = true; // avoid any unwanted network call during capability detection
common_params_handle_models(params, LLAMA_EXAMPLE_SERVER);
if (params.mmproj.path.empty()) {
multimodal = { false, false };
} else {
multimodal = mtmd_get_cap_from_file(params.mmproj.path.c_str());
}
} catch (const std::exception & e) {
LOG_WRN("failed to initialize common_params for multimodal capability detection: %s\n", e.what());
multimodal = { false, false };
}
}
//
// server_models
//
@ -236,6 +260,7 @@ void server_models::add_model(server_model_meta && meta) {
}
meta.update_args(ctx_preset, bin_path); // render args
meta.update_caps();
std::string name = meta.name;
mapping[name] = instance_t{
/* subproc */ std::make_shared<subprocess_s>(),
@ -346,8 +371,10 @@ void server_models::load_models() {
/* status */ SERVER_MODEL_STATUS_UNLOADED,
/* last_used */ 0,
/* args */ std::vector<std::string>(),
/* loaded_info */ {},
/* exit_code */ 0,
/* stop_timeout */ DEFAULT_STOP_TIMEOUT,
/* multimodal */ mtmd_caps{false, false},
};
add_model(std::move(meta));
}
@ -481,6 +508,7 @@ void server_models::load_models() {
inst.meta.exit_code = 0; // clear failed state so the model can be reloaded
inst.meta.update_args(ctx_preset, bin_path);
inst.meta.update_caps();
}
// add models that are new in this reload
@ -496,8 +524,10 @@ void server_models::load_models() {
/* status */ SERVER_MODEL_STATUS_UNLOADED,
/* last_used */ 0,
/* args */ std::vector<std::string>(),
/* loaded_info */ {},
/* exit_code */ 0,
/* stop_timeout */ DEFAULT_STOP_TIMEOUT,
/* multimodal */ mtmd_caps{false, false},
};
add_model(std::move(meta));
newly_added.push_back(name);
@ -1206,14 +1236,28 @@ void server_models_routes::init_routes() {
status["failed"] = true;
}
// pi coding agent multimodal compatibility
json input_modalities = json::array({"text"});
if (meta.multimodal.inp_vision) {
input_modalities.push_back("image");
}
if (meta.multimodal.inp_audio) {
input_modalities.push_back("audio");
}
json architecture {
{"input_modalities", input_modalities},
{"output_modalities", json::array({"text"})},
};
json model_info = json {
{"id", meta.name},
{"aliases", meta.aliases},
{"tags", meta.tags},
{"object", "model"}, // for OAI-compat
{"owned_by", "llamacpp"}, // for OAI-compat
{"created", t}, // for OAI-compat
{"status", status},
{"id", meta.name},
{"aliases", meta.aliases},
{"tags", meta.tags},
{"object", "model"}, // for OAI-compat
{"owned_by", "llamacpp"}, // for OAI-compat
{"created", t}, // for OAI-compat
{"status", status},
{"architecture", architecture},
// TODO: add other fields, may require reading GGUF metadata
};

2
tools/server/server-models.h
View file

@ -66,6 +66,7 @@ struct server_model_meta {
json loaded_info; // info to be reflected via /v1/models endpoint
int exit_code = 0; // exit code of the model instance process (only valid if status == FAILED)
int stop_timeout = 0; // seconds to wait before force-killing the model instance during shutdown
mtmd_caps multimodal; // multimodal capabilities
bool is_ready() const {
return status == SERVER_MODEL_STATUS_LOADED;
@ -80,6 +81,7 @@ struct server_model_meta {
}
void update_args(common_preset_context & ctx_presets, std::string bin_path);
void update_caps();
};
struct subprocess_s;

12
tools/server/server-task.cpp
View file

@ -1988,7 +1988,7 @@ server_prompt * server_prompt_cache::alloc(const server_prompt & prompt, size_t
const int cur_lcp_len = it->tokens.get_common_prefix(prompt.tokens);
if (cur_lcp_len == (int) prompt.tokens.size()) {
SRV_WRN("%s", " - prompt is already in the cache, skipping\n");
SRV_INF("%s", " - prompt is already in the cache, skipping\n");
return nullptr;
}
}
@ -2043,7 +2043,7 @@ bool server_prompt_cache::load(server_prompt & prompt, const server_tokens & tok
float f_keep_best = prompt.tokens.size() > 0 ? float(lcp_best) / prompt.tokens.size() : -1.0f; // empty slot: any cache entry wins
float sim_best = float(lcp_best) / tokens_new.size();
SRV_WRN(" - looking for better prompt, base f_keep = %.3f, sim = %.3f\n", f_keep_best, sim_best);
SRV_INF(" - looking for better prompt, base f_keep = %.3f, sim = %.3f\n", f_keep_best, sim_best);
auto it_best = states.end();
@ -2068,7 +2068,7 @@ bool server_prompt_cache::load(server_prompt & prompt, const server_tokens & tok
}
if (it_best != states.end()) {
SRV_WRN(" - found better prompt with f_keep = %.3f, sim = %.3f\n", f_keep_best, sim_best);
SRV_INF(" - found better prompt with f_keep = %.3f, sim = %.3f\n", f_keep_best, sim_best);
{
auto & data = it_best->data.main;
@ -2076,7 +2076,7 @@ bool server_prompt_cache::load(server_prompt & prompt, const server_tokens & tok
const size_t size = data.size();
const size_t n = llama_state_seq_set_data_ext(ctx_tgt, data.data(), size, id_slot, 0);
if (n != size) {
SRV_WRN("failed to restore state with size %zu\n", size);
SRV_ERR("failed to restore state with size %zu\n", size);
return false;
}
@ -2145,11 +2145,11 @@ void server_prompt_cache::update() {
}
}
SRV_WRN(" - cache state: %zu prompts, %.3f MiB (limits: %.3f MiB, %zu tokens, %zu est)\n",
SRV_INF(" - cache state: %zu prompts, %.3f MiB (limits: %.3f MiB, %zu tokens, %zu est)\n",
states.size(), size() / (1024.0 * 1024.0), limit_size / (1024.0 * 1024.0), limit_tokens, limit_tokens_cur);
for (const auto & state : states) {
SRV_WRN(" - prompt %p: %7d tokens, checkpoints: %2zu, %9.3f MiB\n",
SRV_INF(" - prompt %p: %7d tokens, checkpoints: %2zu, %9.3f MiB\n",
(const void *)&state, state.n_tokens(), state.checkpoints.size(), state.size() / (1024.0 * 1024.0));
}
}

44
tools/server/server.cpp
View file

@ -83,17 +83,22 @@ int main(int argc, char ** argv) {
return 1;
}
llama_backend_init();
llama_numa_init(params.numa);
common_params_print_info(params);
// validate batch size for embeddings
// embeddings require all tokens to be processed in a single ubatch
// see https://github.com/ggml-org/llama.cpp/issues/12836
if (params.embedding && params.n_batch > params.n_ubatch) {
LOG_WRN("%s: embeddings enabled with n_batch (%d) > n_ubatch (%d)\n", __func__, params.n_batch, params.n_ubatch);
LOG_WRN("%s: setting n_batch = n_ubatch = %d to avoid assertion failure\n", __func__, params.n_ubatch);
SRV_WRN("embeddings enabled with n_batch (%d) > n_ubatch (%d)\n", params.n_batch, params.n_ubatch);
SRV_WRN("setting n_batch = n_ubatch = %d to avoid assertion failure\n", params.n_ubatch);
params.n_batch = params.n_ubatch;
}
if (params.n_parallel < 0) {
LOG_INF("%s: n_parallel is set to auto, using n_parallel = 4 and kv_unified = true\n", __func__);
SRV_INF("%s", "n_parallel is set to auto, using n_parallel = 4 and kv_unified = true\n");
params.n_parallel = 4;
params.kv_unified = true;
@ -107,15 +112,9 @@ int main(int argc, char ** argv) {
// struct that contains llama context and inference
server_context ctx_server;
llama_backend_init();
llama_numa_init(params.numa);
LOG_INF("build_info: %s\n", llama_build_info());
LOG_INF("%s\n", common_params_get_system_info(params).c_str());
server_http_context ctx_http;
if (!ctx_http.init(params)) {
LOG_ERR("%s: failed to initialize HTTP server\n", __func__);
SRV_ERR("%s", "failed to initialize HTTP server\n");
return 1;
}
@ -134,7 +133,7 @@ int main(int argc, char ** argv) {
try {
models_routes.emplace(params, argc, argv);
} catch (const std::exception & e) {
LOG_ERR("%s: failed to initialize router models: %s\n", __func__, e.what());
SRV_ERR("failed to initialize router models: %s\n", e.what());
return 1;
}
@ -222,7 +221,7 @@ int main(int argc, char ** argv) {
try {
tools.setup(params.server_tools);
} catch (const std::exception & e) {
LOG_ERR("%s: tools setup failed: %s\n", __func__, e.what());
SRV_ERR("tools setup failed: %s\n", e.what());
return 1;
}
SRV_WRN("%s", "-----------------\n");
@ -240,7 +239,7 @@ int main(int argc, char ** argv) {
std::function<void()> clean_up;
if (is_router_server) {
LOG_INF("%s: starting router server, no model will be loaded in this process\n", __func__);
SRV_INF("%s", "starting router server, no model will be loaded in this process\n");
clean_up = [&models_routes]() {
SRV_INF("%s: cleaning up before exit...\n", __func__);
@ -252,7 +251,7 @@ int main(int argc, char ** argv) {
if (!ctx_http.start()) {
clean_up();
LOG_ERR("%s: exiting due to HTTP server error\n", __func__);
SRV_ERR("%s", "exiting due to HTTP server error\n");
return 1;
}
ctx_http.is_ready.store(true);
@ -273,12 +272,12 @@ int main(int argc, char ** argv) {
// start the HTTP server before loading the model to be able to serve /health requests
if (!ctx_http.start()) {
clean_up();
LOG_ERR("%s: exiting due to HTTP server error\n", __func__);
SRV_ERR("%s", "exiting due to HTTP server error\n");
return 1;
}
// load the model
LOG_INF("%s: loading model\n", __func__);
SRV_INF("%s", "loading model\n");
if (server_models::is_child_server()) {
ctx_server.on_sleeping_changed([&](bool sleeping) {
@ -291,14 +290,14 @@ int main(int argc, char ** argv) {
if (ctx_http.thread.joinable()) {
ctx_http.thread.join();
}
LOG_ERR("%s: exiting due to model loading error\n", __func__);
SRV_ERR("%s", "exiting due to model loading error\n");
return 1;
}
routes.update_meta(ctx_server);
ctx_http.is_ready.store(true);
LOG_INF("%s: model loaded\n", __func__);
SRV_INF("%s", "model loaded\n");
shutdown_handler = [&](int) {
// this will unblock start_loop()
@ -322,9 +321,9 @@ int main(int argc, char ** argv) {
#endif
if (is_router_server) {
LOG_INF("%s: router server is listening on %s\n", __func__, ctx_http.listening_address.c_str());
LOG_INF("%s: NOTE: router mode is experimental\n", __func__);
LOG_INF("%s: it is not recommended to use this mode in untrusted environments\n", __func__);
SRV_INF("router server is listening on %s\n", ctx_http.listening_address.c_str());
SRV_WRN("%s", "NOTE: router mode is experimental\n");
SRV_WRN("%s", " it is not recommended to use this mode in untrusted environments\n");
if (ctx_http.thread.joinable()) {
ctx_http.thread.join(); // keep the main thread alive
}
@ -332,8 +331,7 @@ int main(int argc, char ** argv) {
// when the HTTP server stops, clean up and exit
clean_up();
} else {
LOG_INF("%s: server is listening on %s\n", __func__, ctx_http.listening_address.c_str());
LOG_INF("%s: starting the main loop...\n", __func__);
SRV_INF("server is listening on %s\n", ctx_http.listening_address.c_str());
// optionally, notify router server that this instance is ready
std::thread monitor_thread;

39
tools/server/tests/unit/test_chat_completion.py
View file

@ -178,6 +178,45 @@ def test_chat_template_assistant_prefill(prefill, re_prefill):
assert res.body["__verbose"]["prompt"] == f"<s> <|start_header_id|>system<|end_header_id|>\n\nBook<|eot_id|><|start_header_id|>user<|end_header_id|>\n\nWhat is the best book<|eot_id|><|start_header_id|>assistant<|end_header_id|>\n\n{re_prefill}"
def test_chat_template_continue_final_message_vllm_compat():
"""continue_final_message is the vLLM/transformers explicit alias for the prefill_assistant heuristic.
Both must produce the same prompt."""
global server
server.chat_template = "llama3"
server.debug = True
server.start()
res = server.make_request("POST", "/chat/completions", data={
"max_tokens": 8,
"add_generation_prompt": False,
"continue_final_message": True,
"messages": [
{"role": "system", "content": "Book"},
{"role": "user", "content": "What is the best book"},
{"role": "assistant", "content": "Whill"},
]
})
assert res.status_code == 200
assert "__verbose" in res.body
assert res.body["__verbose"]["prompt"] == "<s> <|start_header_id|>system<|end_header_id|>\n\nBook<|eot_id|><|start_header_id|>user<|end_header_id|>\n\nWhat is the best book<|eot_id|><|start_header_id|>assistant<|end_header_id|>\n\nWhill"
def test_chat_template_continue_final_message_mutual_exclusion():
"""add_generation_prompt and continue_final_message both set to true must be rejected"""
global server
server.chat_template = "llama3"
server.start()
res = server.make_request("POST", "/chat/completions", data={
"max_tokens": 8,
"add_generation_prompt": True,
"continue_final_message": True,
"messages": [
{"role": "user", "content": "Hi"},
{"role": "assistant", "content": "Hello"},
]
})
assert res.status_code == 400
def test_apply_chat_template():
global server
server.chat_template = "command-r"

2
tools/server/webui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessage.svelte
View file

@ -179,7 +179,7 @@
isEditing = false;
// If canceling a new system message with placeholder content, remove it without deleting children
if (message.role === MessageRole.SYSTEM) {
if (message.role === MessageRole.SYSTEM && message.content === SYSTEM_MESSAGE_PLACEHOLDER) {
const conversationDeleted = await chatStore.removeSystemPromptPlaceholder(message.id);
if (conversationDeleted) {

3
tools/server/webui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistant.svelte
View file

@ -74,7 +74,6 @@
const editCtx = getMessageEditContext();
const isAgentic = $derived(hasAgenticContent(message, toolMessages));
const hasReasoning = $derived(!!message.reasoningContent);
const processingState = useProcessingState();
let currentConfig = $derived(config());
@ -329,7 +328,7 @@
{onCopy}
{onEdit}
{onRegenerate}
onContinue={currentConfig.enableContinueGeneration && !hasReasoning ? onContinue : undefined}
onContinue={currentConfig.enableContinueGeneration ? onContinue : undefined}
{onForkConversation}
{onDelete}
{onConfirmDelete}

124
tools/server/webui/src/lib/components/app/chat/ChatMessages/ChatMessages.svelte
View file

@ -1,4 +1,6 @@
<script lang="ts">
import { onMount } from 'svelte';
import { beforeNavigate, afterNavigate } from '$app/navigation';
import { ChatMessage, ChatMessageUserPending } from '$lib/components/app';
import { setChatActionsContext } from '$lib/contexts';
import { MessageRole } from '$lib/enums';
@ -27,11 +29,14 @@
interface Props {
messages?: DatabaseMessage[];
onUserAction?: () => void;
onMessagesReady?: (messageCount: number) => void;
}
let { messages = [], onUserAction }: Props = $props();
let { messages = [], onUserAction, onMessagesReady }: Props = $props();
let allConversationMessages = $state<DatabaseMessage[]>([]);
let isVisible = $state(false);
let previousConversationId = $state<string | null>(null);
const currentConfig = config();
@ -117,15 +122,51 @@
}
}
// Single effect that tracks both conversation and message changes
// Track conversation changes to trigger transition even on same route
$effect(() => {
const conversation = activeConversation();
const currentId = conversation?.id ?? null;
if (conversation) {
refreshAllMessages();
if (currentId !== previousConversationId && previousConversationId !== null) {
// Conversation changed - trigger fade out/in
isVisible = false;
requestAnimationFrame(() => {
refreshAllMessages();
previousConversationId = currentId;
requestAnimationFrame(() => {
isVisible = true;
});
});
} else {
previousConversationId = currentId;
if (conversation) {
refreshAllMessages();
}
}
});
$effect(() => {
void allConversationMessages;
onMessagesReady?.(displayMessages.length);
});
onMount(() => {
requestAnimationFrame(() => {
isVisible = true;
});
});
beforeNavigate(() => {
isVisible = false;
});
afterNavigate(() => {
requestAnimationFrame(() => {
isVisible = true;
});
});
let displayMessages = $derived.by(() => {
if (!messages.length) {
return [];
@ -207,42 +248,47 @@
});
</script>
{#each displayMessages as { message, toolMessages, isLastAssistantMessage, siblingInfo } (message.id)}
<ChatMessage
class="mx-auto mt-12 w-full max-w-[48rem]"
{message}
{toolMessages}
{isLastAssistantMessage}
{siblingInfo}
/>
{/each}
{#if activeConversation() && agenticPendingSteeringMessageContent(activeConversation()!.id)}
{@const convId = activeConversation()!.id}
{@const pendingContent = agenticPendingSteeringMessageContent(convId)}
{#if pendingContent}
<ChatMessageUserPending
<div
class="transition-opacity delay-300 duration-500 ease-out
{isVisible ? 'opacity-100' : 'opacity-0'}"
>
{#each displayMessages as { message, toolMessages, isLastAssistantMessage, siblingInfo } (message.id)}
<ChatMessage
class="mx-auto mt-12 w-full max-w-[48rem]"
content={pendingContent}
extras={agenticPendingSteeringMessageExtras(convId)}
onSendImmediately={() => chatStore.abortCurrentFlow(convId)}
onEdit={(newContent, extras) => agenticInjectSteeringMessage(convId, newContent, extras)}
onDelete={() => agenticClearSteeringMessage(convId)}
{message}
{toolMessages}
{isLastAssistantMessage}
{siblingInfo}
/>
{/if}
{:else if activeConversation() && chatPendingMessageContent(activeConversation()!.id)}
{@const convId = activeConversation()!.id}
{@const pendingContent = chatPendingMessageContent(convId)}
{/each}
{#if pendingContent}
<ChatMessageUserPending
class="mx-auto mt-12 w-full max-w-[48rem]"
content={pendingContent}
extras={chatPendingMessageExtras(convId)}
onSendImmediately={() => chatStore.abortCurrentFlow(convId)}
onEdit={(newContent, extras) => chatInjectPendingMessage(convId, newContent, extras)}
onDelete={() => chatClearPendingMessage(convId)}
/>
{#if activeConversation() && agenticPendingSteeringMessageContent(activeConversation()!.id)}
{@const convId = activeConversation()!.id}
{@const pendingContent = agenticPendingSteeringMessageContent(convId)}
{#if pendingContent}
<ChatMessageUserPending
class="mx-auto mt-12 w-full max-w-[48rem]"
content={pendingContent}
extras={agenticPendingSteeringMessageExtras(convId)}
onSendImmediately={() => chatStore.abortCurrentFlow(convId)}
onEdit={(newContent, extras) => agenticInjectSteeringMessage(convId, newContent, extras)}
onDelete={() => agenticClearSteeringMessage(convId)}
/>
{/if}
{:else if activeConversation() && chatPendingMessageContent(activeConversation()!.id)}
{@const convId = activeConversation()!.id}
{@const pendingContent = chatPendingMessageContent(convId)}
{#if pendingContent}
<ChatMessageUserPending
class="mx-auto mt-12 w-full max-w-[48rem]"
content={pendingContent}
extras={chatPendingMessageExtras(convId)}
onSendImmediately={() => chatStore.abortCurrentFlow(convId)}
onEdit={(newContent, extras) => chatInjectPendingMessage(convId, newContent, extras)}
onDelete={() => chatClearPendingMessage(convId)}
/>
{/if}
{/if}
{/if}
</div>

19
tools/server/webui/src/lib/components/app/chat/ChatScreen/ChatScreen.svelte
View file

@ -42,6 +42,8 @@
let { showCenteredEmpty = false } = $props();
const autoScroll = createAutoScrollController();
let disableAutoScroll = $derived(Boolean(config().disableAutoScroll));
let chatScrollContainer: HTMLDivElement | undefined = $state();
let dragCounter = $state(0);
@ -49,8 +51,6 @@
let showFileErrorDialog = $state(false);
let uploadedFiles = $state<ChatUploadedFile[]>([]);
const autoScroll = createAutoScrollController({ isColumnReverse: true });
let fileErrorData = $state<{
generallyUnsupported: File[];
modalityUnsupported: File[];
@ -322,6 +322,14 @@
}
});
function handleMessagesReady() {
if (!disableAutoScroll && !autoScroll.userScrolledUp) {
requestAnimationFrame(() => {
autoScroll.scrollToBottom('instant');
});
}
}
onMount(() => {
autoScroll.startObserving();
@ -357,7 +365,7 @@
<div
bind:this={chatScrollContainer}
aria-label="Chat interface with file drop zone"
class="flex h-full flex-col-reverse overflow-y-auto px-4 md:px-6"
class="flex h-full flex-col overflow-y-auto px-4 md:px-6"
ondragenter={handleDragEnter}
ondragleave={handleDragLeave}
ondragover={handleDragOver}
@ -371,8 +379,11 @@
messages={activeMessages()}
onUserAction={() => {
autoScroll.enable();
autoScroll.scrollToBottom();
if (!autoScroll.userScrolledUp) {
autoScroll.scrollToBottom();
}
}}
onMessagesReady={handleMessagesReady}
/>
{/if}

21
tools/server/webui/src/lib/components/app/models/ModelId.svelte
View file

@ -31,9 +31,12 @@
let parsed = $derived(ModelsService.parseModelId(modelId));
let resolvedShowRaw = $derived(showRaw ?? (config().showRawModelNames as boolean) ?? false);
let displayName = $derived(parsed.modelName ?? modelId);
let allAliases = $derived(aliases ?? []);
let allTags = $derived([...(parsed.tags ?? []), ...(tags ?? [])]);
let uniqueAliases = $derived([...new Set(aliases ?? [])]);
let uniqueTags = $derived([...new Set([...(parsed.tags ?? []), ...(tags ?? [])])]);
let primaryAlias = $derived(uniqueAliases.length === 1 ? uniqueAliases[0] : null);
let displayName = $derived(primaryAlias ?? parsed.modelName ?? modelId);
</script>
{#if resolvedShowRaw}
@ -56,14 +59,18 @@
</span>
{/if}
{#if allAliases.length > 0}
{#each allAliases as alias (alias)}
{#if primaryAlias}
{#if primaryAlias !== parsed.modelName}
<span class={badgeClass}>{parsed.modelName ?? modelId}</span>
{/if}
{:else if uniqueAliases.length > 1}
{#each uniqueAliases as alias (alias)}
<span class={badgeClass}>{alias}</span>
{/each}
{/if}
{#if allTags.length > 0}
{#each allTags as tag (tag)}
{#if uniqueTags.length > 0}
{#each uniqueTags as tag (tag)}
<span class={tagBadgeClass}>{tag}</span>
{/each}
{/if}

2
tools/server/webui/src/lib/constants/settings-registry.ts
View file

@ -122,7 +122,7 @@ const SETTINGS_REGISTRY: Record<string, SettingsSectionEntry> = {
{
key: SETTINGS_KEYS.ENABLE_CONTINUE_GENERATION,
label: 'Enable "Continue" button',
help: 'Enable "Continue" button for assistant messages. Currently works only with non-reasoning models.',
help: 'Enable "Continue" button for assistant messages, including reasoning models.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.GENERAL,

56
tools/server/webui/src/lib/hooks/use-auto-scroll.svelte.ts
View file

@ -2,7 +2,6 @@ import { AUTO_SCROLL_AT_BOTTOM_THRESHOLD, AUTO_SCROLL_INTERVAL } from '$lib/cons
export interface AutoScrollOptions {
disabled?: boolean;
isColumnReverse?: boolean;
}
/**
@ -12,24 +11,19 @@ export interface AutoScrollOptions {
* - Auto-scrolls to bottom during streaming/loading
* - Stops auto-scroll when user manually scrolls up
* - Resumes auto-scroll when user scrolls back to bottom
* - Supports both normal and column-reverse scroll containers
*/
export class AutoScrollController {
private _autoScrollEnabled = $state(true);
private _userScrolledUp = $state(false);
private _lastScrollTop = $state(0);
private _scrollInterval: ReturnType<typeof setInterval> | undefined;
private _scrollTimeout: ReturnType<typeof setTimeout> | undefined;
private _container: HTMLElement | undefined;
private _disabled: boolean;
private _isColumnReverse: boolean;
private _mutationObserver: MutationObserver | null = null;
private _rafPending = false;
private _observerEnabled = false;
constructor(options: AutoScrollOptions = {}) {
this._disabled = options.disabled ?? false;
this._isColumnReverse = options.isColumnReverse ?? false;
}
get autoScrollEnabled(): boolean {
@ -56,6 +50,7 @@ export class AutoScrollController {
* Updates the disabled state.
*/
setDisabled(disabled: boolean): void {
if (this._disabled === disabled) return;
this._disabled = disabled;
if (disabled) {
this._autoScrollEnabled = false;
@ -73,20 +68,8 @@ export class AutoScrollController {
if (this._disabled || !this._container) return;
const { scrollTop, scrollHeight, clientHeight } = this._container;
let distanceFromBottom: number;
let isScrollingUp: boolean;
if (this._isColumnReverse) {
// column-reverse: scrollTop=0 at bottom, negative when scrolled up
distanceFromBottom = Math.abs(scrollTop);
isScrollingUp = scrollTop < this._lastScrollTop;
} else {
// normal: scrollTop=0 at top, increases when scrolled down
distanceFromBottom = scrollHeight - clientHeight - scrollTop;
isScrollingUp = scrollTop < this._lastScrollTop;
}
const distanceFromBottom = scrollHeight - clientHeight - scrollTop;
const isScrollingUp = scrollTop < this._lastScrollTop;
const isAtBottom = distanceFromBottom < AUTO_SCROLL_AT_BOTTOM_THRESHOLD;
if (isScrollingUp && !isAtBottom) {
@ -97,17 +80,6 @@ export class AutoScrollController {
this._autoScrollEnabled = true;
}
if (this._scrollTimeout) {
clearTimeout(this._scrollTimeout);
}
this._scrollTimeout = setTimeout(() => {
if (isAtBottom) {
this._userScrolledUp = false;
this._autoScrollEnabled = true;
}
}, AUTO_SCROLL_INTERVAL);
this._lastScrollTop = scrollTop;
}
@ -116,13 +88,7 @@ export class AutoScrollController {
*/
scrollToBottom(behavior: ScrollBehavior = 'smooth'): void {
if (this._disabled || !this._container) return;
if (this._isColumnReverse) {
// column-reverse: scrollTop=0 is the bottom
this._container.scrollTo({ top: 0, behavior });
} else {
this._container.scrollTo({ top: this._container.scrollHeight, behavior });
}
this._container.scrollTo({ top: this._container.scrollHeight, behavior });
}
/**
@ -180,11 +146,6 @@ export class AutoScrollController {
destroy(): void {
this.stopInterval();
this._doStopObserving();
if (this._scrollTimeout) {
clearTimeout(this._scrollTimeout);
this._scrollTimeout = undefined;
}
}
/**
@ -210,20 +171,13 @@ export class AutoScrollController {
private _doStartObserving(): void {
if (!this._container || this._mutationObserver) return;
const isReverse = this._isColumnReverse;
this._mutationObserver = new MutationObserver(() => {
if (!this._autoScrollEnabled || this._rafPending) return;
this._rafPending = true;
requestAnimationFrame(() => {
this._rafPending = false;
if (this._autoScrollEnabled && this._container) {
if (isReverse) {
// column-reverse: scrollTop=0 is the bottom
this._container.scrollTop = 0;
} else {
this._container.scrollTop = this._container.scrollHeight;
}
this._container.scrollTop = this._container.scrollHeight;
}
});
});

8
tools/server/webui/src/lib/services/chat.service.ts
View file

@ -130,7 +130,8 @@ export class ChatService {
timings_per_token,
// Config options
disableReasoningParsing,
excludeReasoningFromContext
excludeReasoningFromContext,
continueFinalMessage
} = options;
const normalizedMessages: ApiChatMessageData[] = messages
@ -209,6 +210,11 @@ export class ChatService {
? ReasoningFormat.NONE
: ReasoningFormat.AUTO;
if (continueFinalMessage) {
requestBody.continue_final_message = true;
requestBody.add_generation_prompt = false;
}
if (temperature !== undefined) requestBody.temperature = temperature;
if (max_tokens !== undefined) {
// Set max_tokens to -1 (infinite) when explicitly configured as 0 or null

75
tools/server/webui/src/lib/stores/chat.svelte.ts
View file

@ -674,7 +674,8 @@ class ChatStore {
},
onReasoningChunk: (chunk: string) => {
streamedReasoningContent += chunk;
// Update UI to show reasoning is being received
// mark streaming state so a stop mid-thinking can persist the partial reasoning
this.setChatStreaming(convId, streamedContent, currentMessageId);
const idx = conversationsStore.findMessageIndex(currentMessageId);
conversationsStore.updateMessageAtIndex(idx, {
reasoningContent: streamedReasoningContent
@ -989,38 +990,51 @@ class ChatStore {
const conversationId = convId || conversationsStore.activeConversation?.id;
if (!conversationId) return;
const streamingState = this.getChatStreaming(conversationId);
if (!streamingState || !streamingState.response.trim()) return;
if (!streamingState) return;
const messages =
conversationId === conversationsStore.activeConversation?.id
? conversationsStore.activeMessages
: await conversationsStore.getConversationMessages(conversationId);
if (!messages.length) return;
const lastMessage = messages[messages.length - 1];
if (lastMessage?.role === MessageRole.ASSISTANT) {
try {
const updateData: { content: string; timings?: ChatMessageTimings } = {
content: streamingState.response
};
const lastKnownState = this.getProcessingState(conversationId);
if (lastKnownState) {
updateData.timings = {
prompt_n: lastKnownState.promptTokens || 0,
prompt_ms: lastKnownState.promptMs,
predicted_n: lastKnownState.tokensDecoded || 0,
cache_n: lastKnownState.cacheTokens || 0,
predicted_ms:
lastKnownState.tokensPerSecond && lastKnownState.tokensDecoded
? (lastKnownState.tokensDecoded / lastKnownState.tokensPerSecond) * 1000
: undefined
};
}
await DatabaseService.updateMessage(lastMessage.id, updateData);
lastMessage.content = streamingState.response;
if (updateData.timings) lastMessage.timings = updateData.timings;
} catch (error) {
lastMessage.content = streamingState.response;
console.error('Failed to save partial response:', error);
if (lastMessage?.role !== MessageRole.ASSISTANT) return;
const partialContent = streamingState.response;
const partialReasoning = lastMessage.reasoningContent || '';
// nothing to persist when both content and reasoning are empty (e.g. stop before any token)
if (!partialContent.trim() && !partialReasoning.trim()) return;
try {
const updateData: {
content: string;
reasoningContent?: string;
timings?: ChatMessageTimings;
} = {
content: partialContent
};
if (partialReasoning) {
updateData.reasoningContent = partialReasoning;
}
const lastKnownState = this.getProcessingState(conversationId);
if (lastKnownState) {
updateData.timings = {
prompt_n: lastKnownState.promptTokens || 0,
prompt_ms: lastKnownState.promptMs,
predicted_n: lastKnownState.tokensDecoded || 0,
cache_n: lastKnownState.cacheTokens || 0,
predicted_ms:
lastKnownState.tokensPerSecond && lastKnownState.tokensDecoded
? (lastKnownState.tokensDecoded / lastKnownState.tokensPerSecond) * 1000
: undefined
};
}
await DatabaseService.updateMessage(lastMessage.id, updateData);
lastMessage.content = partialContent;
if (updateData.timings) lastMessage.timings = updateData.timings;
} catch (error) {
lastMessage.content = partialContent;
console.error('Failed to save partial response:', error);
}
}
@ -1265,7 +1279,11 @@ class ChatStore {
const conversationContext = conversationsStore.activeMessages.slice(0, idx);
const contextWithContinue = [
...conversationContext,
{ role: MessageRole.ASSISTANT as const, content: originalContent }
{
role: MessageRole.ASSISTANT as const,
content: originalContent,
reasoning_content: originalReasoning || undefined
}
];
let appendedContent = '';
@ -1283,6 +1301,7 @@ class ChatStore {
contextWithContinue,
{
...this.getApiOptions(),
continueFinalMessage: true,
onChunk: (chunk: string) => {
appendedContent += chunk;
hasReceivedContent = true;
@ -1291,6 +1310,8 @@ class ChatStore {
onReasoningChunk: (chunk: string) => {
appendedReasoning += chunk;
hasReceivedContent = true;
// mark streaming state so a stop mid-thinking can persist the partial reasoning
this.setChatStreaming(msg.convId, originalContent + appendedContent, msg.id);
conversationsStore.updateMessageAtIndex(idx, {
reasoningContent: originalReasoning + appendedReasoning
});

3
tools/server/webui/src/lib/types/api.d.ts vendored
View file

@ -239,6 +239,9 @@ export interface ApiChatCompletionRequest {
// Custom parameters (JSON string)
custom?: Record<string, unknown>;
timings_per_token?: boolean;
// Continuation control (vLLM compat)
add_generation_prompt?: boolean;
continue_final_message?: boolean;
}
export interface ApiChatCompletionToolCallFunctionDelta {

2
tools/server/webui/src/lib/types/settings.d.ts vendored
View file

@ -92,6 +92,8 @@ export interface SettingsChatServiceOptions {
// Custom parameters
custom?: string;
timings_per_token?: boolean;
// Continuation control (vLLM compat), opt in to the explicit continue final message flag
continueFinalMessage?: boolean;
// Callbacks
onChunk?: (chunk: string) => void;
onReasoningChunk?: (chunk: string) => void;

2
tools/tts/tts.cpp
View file

@ -604,8 +604,8 @@ int main(int argc, char ** argv) {
}
LOG_INF("sampler seed: %u\n", common_sampler_get_seed(smpl[0]));
LOG_INF("sampler params: \n%s\n", params.sampling.print().c_str());
LOG_INF("sampler chain: %s\n", common_sampler_print(smpl[0]).c_str());
LOG_INF("sampler params: \n%s\n", params.sampling.print().c_str());
LOG_INF("%s: loading done\n", __func__);