Compare commits

..

26 Commits

Author SHA1 Message Date
Aman Gupta 33a75f41c3 DeepseekV4: reduce graph splits (#25702) 2026-07-15 15:47:18 +08:00
Neo Zhang d3fba0c79d sycl : fix get_rows Q2_K, Q4_K, Q5_K (#25656) 2026-07-15 10:32:28 +03:00
Neo Zhang ae9291e16b sycl : support kernel type fp16 for conv2d_dw (#25653) 2026-07-15 10:31:10 +03:00
Andrew Smith 22b208b1ca sycl : implement xielu op (#25550) 2026-07-15 10:29:12 +03:00
Francois Dugast 0e148a573f sycl: Increase minimum buffer size for USM system allocations (#25525)
Raise the threshold for minimum buffer size from 1 GiB to 4 GiB, based
on real-world experiments of overcommitting device memory with model
weights larger than available VRAM, for example Qwen3.5-35B-A3B-Q8
running on a B70.

Also add a debug message to better track USM system allocations.

Signed-off-by: Francois Dugast <francois.dugast@intel.com>
2026-07-15 10:28:24 +03:00
hmscider 32b741c336 [SYCL] Flash Attention with XMX engine via oneDNN (#25222)
* [SYCL] F16 (default) Flash Attention with XMX engine via oneDNN graph API; Qwen3.6-27b-Q8_0 prefill speed up x1.21 at p=512 and x4.26 at p=80k

* [SYCL] Address review on FA oneDNN path. Result: llama-bench---pp512; 32% increase with fa1; llama-perplexity---0.11% difference; tested model: mradermacher/Meta-Llama-3.1-8B-Instruct-Q8_0.gguf

* PR-25222 revision v2: addressed audits

* [SYCL] flash-attn oneDNN SDPA KV F16 rev 3.0: add BMG gate + multi-device sync. Narrow the scrope of this PR to Battlemage only (bmg; Xe2). Other archs (e.g., alchemist) fall back to existing FA kernel. When device_count >1, apply stream -> wait_and_throw(), validated working path for multi-gpu sync fix by @maxious.

Co-authored-by: maxious <81432+maxious@users.noreply.github.com>

* updated comment on bmg gate, noted the issue

---------

Co-authored-by: scientist3 <scientist.3@users.noreply.github.com>
Co-authored-by: hmscider <hmscider@users.noreply.github.com>
Co-authored-by: maxious <81432+maxious@users.noreply.github.com>
2026-07-15 10:26:53 +03:00
Hongqiang Wang 12127defda opencl: do not use clCreateBufferWithProperties when targeting CL 2.x (#25673) 2026-07-14 19:53:56 -07:00
Hongqiang Wang 00fa7cb284 opencl: handle OOB write in noshuffle GEMV kernels (odd ne01) (#25640) 2026-07-14 13:46:54 -07:00
Hongqiang Wang a4ce2595c5 opencl: avoid the vec path in GEMV for unaligned row stride (#25671)
The f16 GEMV kernels take a vectorized path for ne00 >= 128 that casts the row
pointers to half4 or float4. When the row stride is not aligned, the wide load
becomes misaligned. On devices that require natural alignment for vector loads,
the kernel reads garbage. This is the case Intel GPUs and the kernels produce
incorrect results there. Adreno happpens to be byte addressable and the kernels
happen to work.
2026-07-14 12:27:56 -07:00
Chyan c71854292f hexagon: fix hmx-queue signal enum-narrowing problem (#25677) 2026-07-14 12:27:09 -07:00
Georgi Gerganov bf2c86ddc0 server : refactor prompt cache state ownership (#25649)
* server : clear checkpoints upon prompt clear

* server : move the prompt state data to the server_prompt_cache

Assisted-by: pi:llama.cpp/Qwen3.6-27B

* server : handle batched slot being cleared
2026-07-14 18:25:52 +03:00
Xuan-Son Nguyen 6e52db5b72 server: add --cors-* options (#25655)
* server: add --cors-* options

* add special "localhost" value

* add tests

* fix test

* add link to PR
2026-07-14 17:23:44 +02:00
Bill Sideris 236ab574e0 ui: Fix spacing in tool-call request (#25634) 2026-07-14 17:23:11 +02:00
Emanuil Rusev dfba90db63 webui: parse effective-parameter sizes (E2B, E4B) as params (#25529) 2026-07-14 17:12:22 +02:00
Hongqiang Wang 00e79f6fb1 opencl: fix a dp4a bug for devices where cl_khr_integer_dot_product is unavailable (#25639)
* opencl: do not fail backend init on devices without cl_khr_integer_dot_product

* opencl: do not call dp4 kernels when dp is unavailable

---------

Co-authored-by: Li He <lih@qti.qualcomm.com>
2026-07-14 08:08:13 -07:00
Pascal 17a05e451f ui: fix mcp panel for toggle + timeout + proxy + ON/OFF state (#25631)
* ui: fix MCP panel regressions after settings rework

Restore the llama-server proxy switch in the Add New Server dialog.
The dialog never passed useProxy/onUseProxyChange to McpServerForm,
which only renders the proxy switch when the handler is provided.
The flag is now wired, persisted on addServer, and reset on close.

Bound the MCP connection handshake with the configured timeout.
handshakeTimeoutMs was set in the server config but never consumed.
The SDK timeout only covers the initialize request, not
transport.start(), which can hang forever on an unreachable host.
The whole handshake now races against the timeout and closes the
transport on expiry so the underlying fetch or socket is aborted.

Keep disabled MCP servers visible in management and chat-add UIs.
Collapsing mcpDefaultServerOverrides into mcpServers[i].enabled turned
the visibleMcpServers enabled filter into a visibility trap: toggling
a server off outside a conversation hid it from every surface with no
way to re-enable it. The filter is dropped, tools derived from health
checks still skip disabled servers, and the settings page and server
card render the real card instead of a skeleton for disabled servers
that never receive a startup health check.

* ui: clarify MCP server list semantics and add regression test

Remove the visibleMcpServers getter, a filterless alias of getServers
whose name invites the next refactor to put a filter back. Call sites
read getServers directly, the duplicate list in the chat submenu is
merged, and the misleading local variable in the sheet is renamed.

A parser unit test pins the invariant: enabled is an on/off state,
never a visibility filter, so disabled servers stay listed and
toggleable.

* ui: apply the MCP request timeout setting live to all servers

The per-server requestTimeoutSeconds field was never editable in any
UI and froze the global setting at server creation time, so changing
the timeout in Settings was a no-op for existing servers. The field
is removed from the data model and parsers, the timeout is read live
from the global setting wherever a request config is built, and the
misleading "Can be overridden per server" help text is dropped. A
parser unit test guards against reintroducing the stored field.

* ui: move the MCP request timeout into the Agentic settings section

The MCP section held a single setting. The timeout is a global tool
execution parameter like the other Agentic entries, so it moves there
and the section is removed. Same settings key, no migration needed.

* ui: remove the dead tool preview lines setting

The agenticMaxToolPreviewLines setting was read into AgenticConfig
and consumed by nothing: the agentic loop only uses enabled and
maxTurns. Its help text described a previous architecture where only
truncated previews and the final response survived the loop; tool
results and intermediate turns now persist as full DB messages, so
the setting had no effect at any value. Stale keys in localStorage
or a server ui-config are ignored.

* ui: resolve absent MCP per-chat overrides to the server enabled flag

New conversations started with every MCP server off: the settings
rework stopped seeding a per-conversation override list, assuming
the enabled check would fall back to mcpServers[i].enabled, but it
fell back to false, and the send path passed the raw stored list
with no fallback at all. The per-conversation list is now sparse by
contract, holding only explicit toggles, and every access point
resolves a missing entry to the server's own enabled flag: the
toggle display, the resolved list handed to the agentic flow, and
the enabled check itself.
2026-07-14 16:50:44 +02:00
Aman Gupta 7f575c39d6 DeepseekV4: fix seq_rm (#25588)
* DeepseekV4: fix seq_rm

* implement proper seq_cp

* create actual update context
2026-07-14 21:45:36 +08:00
Jeff Bolz 7cbd61002d vulkan/cpu: Support f16 as SET_ROWS src. (#25432)
* vulkan/cpu: Support f16 as SET_ROWS src.

This adds full support for f16 SET_ROWS (equivalent to f32) to vulkan and CPU
backends, and adds more backend tests.

* Set DenormPreserve 16 when supported, to try to fix failures on Intel

* tune error threshold

* update metal supports_op
2026-07-14 08:26:55 -05:00
Adrien Gallouët 8ff8c4299d tokenize : align usage by using common args (#25516)
Migrate the tokenize tool to common_params_parse, replacing its
hand-rolled argv parsing, Windows UTF-8 handling and file reading
with the shared common helpers.

Expose the model-sourcing flags (-m, -mu, -dr, -hf, -hff, --offline,
HF_TOKEN) to LLAMA_EXAMPLE_TOKENIZE, and register --ids, --stdin,
--no-bos, --no-parse-special and --show-count as common args.
parse_special defaults to true for TOKENIZE to preserve the old
behavior. Errors now go through LOG_ERR instead of fprintf(stderr).

Signed-off-by: Adrien Gallouët <angt@huggingface.co>
2026-07-14 15:20:53 +02:00
fairydreaming a7312ae94f ggml : add a set of functions for checking contiguity of inner tensor dimensions (#25650)
Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com>
2026-07-14 14:37:52 +02:00
Christian Kastner 657e01125a tests: export-graph-ops: exit gracefully when called w/o arguments (#25619)
Fixes a segfault when `test-export-graph-ops` is called without any
arguments.
2026-07-14 13:15:41 +03:00
JusteLeo 47a39665e7 ggml: uniformize im2col dst_type for all conv ops (#23660)
* ggml: uniformize im2col dst_type for all conv ops

* Update ggml/src/ggml.c

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>

* ggml : uniformize im2col casting logic across all conv ops

* fix : allow im2col_f16 to accept any kernel type

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2026-07-14 13:13:13 +03:00
Charles Xu 47c786924a kleidiai : add SME2 f32 kernel (#24414)
* kleidiai : add SME2 f32 kernel

* enable dynamic scheduling for SME2 f32 kernel
2026-07-14 13:12:18 +03:00
Pascal c9330ed0cf ui: add reasoning effort control to mobile add sheet (#25539)
The mobile "+" sheet was missing the reasoning effort section present
in the desktop dropdown, so thinking could not be toggled on touch.

Extract the shared derivation and selection logic into useReasoningMenu
and consume it from both the desktop submenu and the mobile sheet,
keeping a single source of truth and preserving each surface idiom.
2026-07-14 12:05:40 +02:00
Thiago Padilha cb489bc0fb convert_hf_to_gguf: support split MTP export for HY V3 (#25641)
- Add a supports_mtp_export capability to ModelBase so architectures can opt
  into --mtp and --no-mtp without extending a central class allowlist.
- Enable the capability for the existing Qwen3.5/3.6 and Step3.5/3.7
  implementations, and for HY V3, whose converter already supports
  filtering the appended MTP layers.
2026-07-14 11:43:15 +02:00
Christian Kastner ec0dbef816 arg: Flush log before exiting after usage() (#25504)
Under certain conditions, it's possible for messages emitted via LOG()
to get lost before exit, apparently because they are emitted by another
thread. common_params_print_usage() uses printf directly, and is not
affected.

Flushing the log before exit seems to resolve this.
2026-07-14 12:03:22 +03:00
84 changed files with 2111 additions and 881 deletions
+92 -10
View File
@@ -697,7 +697,7 @@ static bool common_params_parse_ex(int argc, char ** argv, common_params_context
}
};
// parse the first time to get -hf option (used for remote preset)
// parse all CLI args now, so that -hf is available below for remote preset resolution
parse_cli_args();
postprocess_cpu_params(params.cpuparams, nullptr);
@@ -748,6 +748,11 @@ static bool common_params_parse_ex(int argc, char ** argv, common_params_context
params.kv_overrides.back().key[0] = 0;
}
if (!params.server_tools.empty() && !params.cors_origins_explicit) {
LOG_WRN("server tools are enabled, using localhost as default CORS origin (change via --cors-origins)\n");
params.cors_origins = "localhost";
}
// pad tensor_buft_overrides for llama_params_fit:
const size_t ntbo = llama_max_tensor_buft_overrides();
while (params.tensor_buft_overrides.size() < ntbo) {
@@ -1077,6 +1082,7 @@ bool common_params_parse(int argc, char ** argv, common_params & params, llama_e
if (ctx_arg.print_usage) {
ctx_arg.print_usage(argc, argv);
}
common_log_flush(common_log_main());
exit(0);
}
if (ctx_arg.params.completion) {
@@ -1178,6 +1184,8 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
params.sampling.temp = 0.2; // lower temp by default for better quality
} else if (ex == LLAMA_EXAMPLE_SERVER) {
params.n_parallel = -1; // auto by default
} else if (ex == LLAMA_EXAMPLE_TOKENIZE) {
params.parse_special = true; // parse special tokens by default, like the old tokenize tool
}
params.use_color = tty_can_use_colors();
@@ -2745,14 +2753,14 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
[](common_params & params, const std::string & value) {
params.model.path = value;
}
).set_examples({LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_EXPORT_LORA, LLAMA_EXAMPLE_DOWNLOAD}).set_env("LLAMA_ARG_MODEL"));
).set_examples({LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_EXPORT_LORA, LLAMA_EXAMPLE_DOWNLOAD, LLAMA_EXAMPLE_TOKENIZE}).set_env("LLAMA_ARG_MODEL"));
add_opt(common_arg(
{"-mu", "--model-url"}, "MODEL_URL",
"model download url (default: unused)",
[](common_params & params, const std::string & value) {
params.model.url = value;
}
).set_examples({LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_DOWNLOAD}).set_env("LLAMA_ARG_MODEL_URL"));
).set_examples({LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_DOWNLOAD, LLAMA_EXAMPLE_TOKENIZE}).set_env("LLAMA_ARG_MODEL_URL"));
add_opt(common_arg(
{ "-dr", "--docker-repo" }, "[<repo>/]<model>[:quant]",
"Docker Hub model repository. repo is optional, default to ai/. quant is optional, default to :latest.\n"
@@ -2761,7 +2769,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
[](common_params & params, const std::string & value) {
params.model.docker_repo = value;
}
).set_examples({LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_DOWNLOAD}).set_env("LLAMA_ARG_DOCKER_REPO"));
).set_examples({LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_DOWNLOAD, LLAMA_EXAMPLE_TOKENIZE}).set_env("LLAMA_ARG_DOCKER_REPO"));
add_opt(common_arg(
{"-hf", "-hfr", "--hf-repo"}, "<user>/<model>[:quant]",
"Hugging Face model repository; quant is optional, case-insensitive, default to Q4_K_M, or falls back to the first file in the repo if Q4_K_M doesn't exist.\n"
@@ -2771,14 +2779,14 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
[](common_params & params, const std::string & value) {
params.model.hf_repo = value;
}
).set_examples({LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_DOWNLOAD}).set_env("LLAMA_ARG_HF_REPO"));
).set_examples({LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_DOWNLOAD, LLAMA_EXAMPLE_TOKENIZE}).set_env("LLAMA_ARG_HF_REPO"));
add_opt(common_arg(
{"-hff", "--hf-file"}, "FILE",
"Hugging Face model file. If specified, it will override the quant in --hf-repo (default: unused)",
[](common_params & params, const std::string & value) {
params.model.hf_file = value;
}
).set_examples({LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_DOWNLOAD}).set_env("LLAMA_ARG_HF_FILE"));
).set_examples({LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_DOWNLOAD, LLAMA_EXAMPLE_TOKENIZE}).set_env("LLAMA_ARG_HF_FILE"));
add_opt(common_arg(
{"-hfv", "-hfrv", "--hf-repo-v"}, "<user>/<model>[:quant]",
"Hugging Face model repository for the vocoder model (default: unused)",
@@ -2799,7 +2807,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
[](common_params & params, const std::string & value) {
params.hf_token = value;
}
).set_examples({LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_DOWNLOAD}).set_env("HF_TOKEN"));
).set_examples({LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_DOWNLOAD, LLAMA_EXAMPLE_TOKENIZE}).set_env("HF_TOKEN"));
add_opt(common_arg(
{"--mtp"},
"also download the multi-token prediction (MTP) head, if available (default: unused)",
@@ -2915,6 +2923,41 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
params.parse_special = true;
}
).set_examples({LLAMA_EXAMPLE_IMATRIX}));
add_opt(common_arg(
{"--ids"},
string_format("only print the token IDs, in a Python-parseable list form like [1, 2, 3] (default: %s)", params.tokenize_ids ? "true" : "false"),
[](common_params & params) {
params.tokenize_ids = true;
}
).set_examples({LLAMA_EXAMPLE_TOKENIZE}));
add_opt(common_arg(
{"--stdin"},
string_format("read the prompt from stdin (mutually exclusive with -f/--file and -p/--prompt) (default: %s)", params.tokenize_stdin ? "true" : "false"),
[](common_params & params) {
params.tokenize_stdin = true;
}
).set_examples({LLAMA_EXAMPLE_TOKENIZE}));
add_opt(common_arg(
{"--no-bos"},
string_format("do not add a BOS token to the prompt, even if the model normally uses one (default: %s)", params.tokenize_no_bos ? "true" : "false"),
[](common_params & params) {
params.tokenize_no_bos = true;
}
).set_examples({LLAMA_EXAMPLE_TOKENIZE}));
add_opt(common_arg(
{"--no-parse-special"},
string_format("do not parse special tokens (chat, tool, etc) (default: %s)", !params.parse_special ? "true" : "false"),
[](common_params & params) {
params.parse_special = false;
}
).set_examples({LLAMA_EXAMPLE_TOKENIZE}));
add_opt(common_arg(
{"--show-count"},
string_format("print the total number of tokens (default: %s)", params.tokenize_show_count ? "true" : "false"),
[](common_params & params) {
params.tokenize_show_count = true;
}
).set_examples({LLAMA_EXAMPLE_TOKENIZE}));
add_opt(common_arg(
{"-pps"},
string_format("is the prompt shared across parallel sequences (default: %s)", params.is_pp_shared ? "true" : "false"),
@@ -3009,6 +3052,42 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
params.public_path = value;
}
).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_STATIC_PATH"));
add_opt(common_arg(
{"--cors-origins"}, "ORIGINS",
string_format(
"comma-separated list of allowed origins for CORS (default: %s)\n"
"if set to special value 'localhost', reflect the Origin header only if it is localhost",
params.cors_origins.c_str()),
[](common_params & params, const std::string & value) {
params.cors_origins = value;
params.cors_origins_explicit = true;
}
).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_CORS_ORIGINS"));
add_opt(common_arg(
{"--cors-methods"}, "METHODS",
string_format("comma-separated list of allowed methods for CORS (default: %s)", params.cors_methods.c_str()),
[](common_params & params, const std::string & value) {
params.cors_methods = value;
}
).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_CORS_METHODS"));
add_opt(common_arg(
{"--cors-headers"}, "HEADERS",
string_format("comma-separated list of allowed headers for CORS (default: %s)", params.cors_headers.c_str()),
[](common_params & params, const std::string & value) {
params.cors_headers = value;
}
).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_CORS_HEADERS"));
add_opt(common_arg(
{"--cors-credentials"},
{"--no-cors-credentials"},
string_format(
"whether to allow credentials for CORS (default: %s)\n"
"note: if this is enabled and --cors-origins is set to * (default), the Origin header will be echoed back, and credentials will always be allowed",
params.cors_credentials ? "enabled" : "disabled"),
[](common_params & params, bool value) {
params.cors_credentials = value;
}
).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_CORS_CREDENTIALS"));
add_opt(common_arg(
{"--api-prefix"}, "PREFIX",
string_format("prefix path the server serves from, without the trailing slash (default: %s)", params.api_prefix.c_str()),
@@ -3042,7 +3121,8 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
{"--tools"}, "TOOL1,TOOL2,...",
"experimental: whether to enable built-in tools for AI agents - do not enable in untrusted environments (default: no tools)\n"
"specify \"all\" to enable all tools\n"
"available tools: read_file, file_glob_search, grep_search, exec_shell_command, write_file, edit_file, get_datetime",
"available tools: read_file, file_glob_search, grep_search, exec_shell_command, write_file, edit_file, get_datetime\n"
"note: for security reasons, this will limit --cors-origins to localhost by default",
[](common_params & params, const std::string & value) {
params.server_tools = parse_csv_row(value);
}
@@ -3050,7 +3130,8 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
add_opt(common_arg(
{"-ag", "--agent"},
{"-no-ag", "--no-agent"},
"whether to enable CORS proxy and all built-in tools - do not enable in untrusted environments (default: disabled)",
"whether to enable CORS proxy and all built-in tools - do not enable in untrusted environments (default: disabled)\n"
"note: for security reasons, this will limit --cors-origins to localhost by default",
[](common_params & params, bool value) {
if (value) {
params.server_tools = {"all"};
@@ -3059,6 +3140,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
params.server_tools.clear();
params.ui_mcp_proxy = false;
}
// note: do not modify cors_origins here, as the options are not evaluated in order (user may explicitly set --cors-origins before --agent)
}
).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_AGENT"));
add_opt(common_arg(
@@ -3505,7 +3587,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
[](common_params & params) {
params.offline = true;
}
).set_examples({LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_DOWNLOAD}).set_env("LLAMA_ARG_OFFLINE"));
).set_examples({LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_DOWNLOAD, LLAMA_EXAMPLE_TOKENIZE}).set_env("LLAMA_ARG_OFFLINE"));
add_opt(common_arg(
{"-lv", "--verbosity", "--log-verbosity"}, "N",
string_format("Set the verbosity threshold. Messages with a higher verbosity will be ignored. Values:\n"
+15
View File
@@ -105,6 +105,7 @@ enum llama_example {
LLAMA_EXAMPLE_RESULTS,
LLAMA_EXAMPLE_EXPORT_GRAPH_OPS,
LLAMA_EXAMPLE_DOWNLOAD,
LLAMA_EXAMPLE_TOKENIZE,
LLAMA_EXAMPLE_COUNT,
};
@@ -630,6 +631,14 @@ struct common_params {
std::string api_prefix = ""; // NOLINT
std::string chat_template = ""; // NOLINT
bool use_jinja = true; // NOLINT
// server CORS params
std::string cors_origins = "*";
std::string cors_methods = "GET, POST, DELETE, OPTIONS";
std::string cors_headers = "*";
bool cors_credentials = true;
bool cors_origins_explicit = false; // for --agent option
bool enable_chat_template = true;
bool force_pure_content_parser = false;
common_reasoning_format reasoning_format = COMMON_REASONING_FORMAT_DEEPSEEK;
@@ -716,6 +725,12 @@ struct common_params {
// batched-bench params
bool batched_bench_output_jsonl = false;
// tokenize params
bool tokenize_ids = false; // if true, only print the token IDs
bool tokenize_stdin = false; // if true, read the prompt from stdin
bool tokenize_no_bos = false; // if true, do not add the BOS token
bool tokenize_show_count = false; // if true, print the total token count
// common params
std::string out_file; // output filename for all example programs
// optional callback for model loading progress and cancellation:
+3 -1
View File
@@ -109,7 +109,9 @@ class ModelBase:
sentence_transformers_dense_modules: bool = False
# MTP (multi-token prediction) export modes; set by main() before instantiation.
# Architectures opt in by overriding the handling (see _Qwen35MtpMixin).
# Architectures that implement the filtering/export behavior opt in by
# setting supports_mtp_export = True on their model class or a mixin.
supports_mtp_export: bool = False
mtp_only: bool = False
no_mtp: bool = False
+1
View File
@@ -361,6 +361,7 @@ class HunyuanVLTextModel(HunYuanModel):
@ModelBase.register("HYV3ForCausalLM")
class HYV3Model(TextModel):
model_arch = gguf.MODEL_ARCH.HY_V3
supports_mtp_export = True
# Trunk layer count, stashed before indexing so the classmethod
# filter_tensors can identify the appended MTP block(s) (mirrors
+1
View File
@@ -541,6 +541,7 @@ class _Qwen35MtpMixin:
`mtp.*` to the standard layer-indexed nextn naming so the existing
tensor_map handles them."""
supports_mtp_export = True
hparams: dict[str, Any]
model_arch: gguf.MODEL_ARCH
gguf_writer: gguf.GGUFWriter
+1
View File
@@ -98,6 +98,7 @@ class Step3VLTextModel(Qwen3Model):
@ModelBase.register("Step3p5ForCausalLM", "Step3p7ForConditionalGeneration")
class Step35Model(TextModel):
model_arch = gguf.MODEL_ARCH.STEP35
supports_mtp_export = True
# The --mtp / --no-mtp toggles are ModelBase.mtp_only / no_mtp (set in
# convert_hf_to_gguf.py main()). Unlike Qwen3.5, which stores MTP under a
+2 -4
View File
@@ -259,10 +259,8 @@ def main() -> None:
sys.exit(1)
if args.mtp or args.no_mtp:
from conversion.qwen import _Qwen35MtpMixin
from conversion.step3 import Step35Model
if not (issubclass(model_class, _Qwen35MtpMixin) or issubclass(model_class, Step35Model)):
logger.error("--mtp / --no-mtp are only supported for Qwen3.5/3.6 and Step3.5 text variants today")
if not model_class.supports_mtp_export:
logger.error("--mtp / --no-mtp are not supported for %s", model_architecture)
sys.exit(1)
if args.no_mtp:
model_class.no_mtp = True
+1 -1
View File
@@ -120,4 +120,4 @@ Legend:
| TRI | ❌ | ❌ | ✅ | ✅ | ✅ | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
| TRUNC | ❌ | ❌ | ✅ | 🟡 | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
| UPSCALE | ❌ | 🟡 | ✅ | ✅ | ❌ | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
| XIELU | ❌ | ❌ | ✅ | ❌ | ❌ | ✅ | ❌ | | ✅ | ✅ | ❌ | ❌ |
| XIELU | ❌ | ❌ | ✅ | ❌ | ❌ | ✅ | ❌ | | ✅ | ✅ | ❌ | ❌ |
+4 -4
View File
@@ -11600,10 +11600,10 @@ zjy 2
"SYCL0","CUMSUM","type=f32,ne=[242004,1,1,1]","support","1","yes","SYCL"
"SYCL0","CUMSUM","type=f32,ne=[375960,1,1,1]","support","1","yes","SYCL"
"SYCL0","CUMSUM","type=f32,ne=[20481,4,1,1]","support","1","yes","SYCL"
"SYCL0","XIELU","type=f32,ne=[10,5,4,3]","support","0","no","SYCL"
"SYCL0","XIELU","type=f16,ne=[10,5,4,3]","support","0","no","SYCL"
"SYCL0","XIELU","type=f32,ne=[512,16,1,1]","support","0","no","SYCL"
"SYCL0","XIELU","type=f16,ne=[512,16,1,1]","support","0","no","SYCL"
"SYCL0","XIELU","type=f32,ne=[10,5,4,3]","support","1","yes","SYCL"
"SYCL0","XIELU","type=f16,ne=[10,5,4,3]","support","1","yes","SYCL"
"SYCL0","XIELU","type=f32,ne=[512,16,1,1]","support","1","yes","SYCL"
"SYCL0","XIELU","type=f16,ne=[512,16,1,1]","support","1","yes","SYCL"
"SYCL0","TRI","type=f32,ne=[10,10,4,3],tri_type=3","support","1","yes","SYCL"
"SYCL0","TRI","type=f32,ne=[10,10,4,3],tri_type=2","support","1","yes","SYCL"
"SYCL0","TRI","type=f32,ne=[10,10,4,3],tri_type=1","support","1","yes","SYCL"
Can't render this file because it is too large.
+4
View File
@@ -780,6 +780,10 @@ extern "C" {
GGML_API bool ggml_is_contiguous_1(const struct ggml_tensor * tensor); // contiguous for dims >= 1
GGML_API bool ggml_is_contiguous_2(const struct ggml_tensor * tensor); // contiguous for dims >= 2
GGML_API bool ggml_is_contiguous_to_1(const struct ggml_tensor * tensor); // contiguous for dims < 1
GGML_API bool ggml_is_contiguous_to_2(const struct ggml_tensor * tensor); // contiguous for dims < 2
GGML_API bool ggml_is_contiguous_to_3(const struct ggml_tensor * tensor); // contiguous for dims < 3
// returns whether the tensor elements are allocated as one contiguous block of memory (no gaps, but permutation ok)
GGML_API bool ggml_is_contiguously_allocated(const struct ggml_tensor * tensor);
+7
View File
@@ -638,6 +638,7 @@ function(ggml_add_cpu_backend_variant_impl tag_name)
${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_f32_qai8dxp_qsi8cxp/
${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_fp32_bf16p_bf16p/
${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_f32_f16p_qsi4c32p/
${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_f32_f32p_f32p/
${KLEIDIAI_SRC}/kai/ukernels/matmul/pack/)
set(ARCH_FLAGS_TEMP "${ARCH_FLAGS}")
@@ -687,9 +688,15 @@ function(ggml_add_cpu_backend_variant_impl tag_name)
${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_fp32_bf16p_bf16p/kai_matmul_clamp_f32_bf16p2vlx2_bf16p2vlx2_2vlx2vl_sme2_mopa_asm.S
${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_f32_f16p_qsi4c32p/kai_matmul_clamp_f32_f16p1vlx2_qsi4c32p4vlx2_1vlx4vl_sme2_mopa.c
${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_f32_f16p_qsi4c32p/kai_matmul_clamp_f32_f16p1vlx2_qsi4c32p4vlx2_1vlx4vl_sme2_mopa_asm.S
${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_f32_f32p_f32p/kai_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa.c
${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_f32_f32p_f32p/kai_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa_asm.S
${KLEIDIAI_SRC}/kai/ukernels/matmul/pack/kai_lhs_pack_bf16p2vlx2_f32_sme.c
${KLEIDIAI_SRC}/kai/ukernels/matmul/pack/kai_rhs_pack_kxn_bf16p2vlx2b_f32_x32_sme.c
${KLEIDIAI_SRC}/kai/ukernels/matmul/pack/kai_lhs_pack_f16pmrx2_f32_neon.c
${KLEIDIAI_SRC}/kai/ukernels/matmul/pack/kai_lhs_pack_f32p2vlx1_f32_sme.c
${KLEIDIAI_SRC}/kai/ukernels/matmul/pack/kai_lhs_pack_f32p2vlx1_f32_sme_asm.S
${KLEIDIAI_SRC}/kai/ukernels/matmul/pack/kai_rhs_pack_nxk_f32p2vlx1biasf32_f32_f32_sme.c
${KLEIDIAI_SRC}/kai/ukernels/matmul/pack/kai_rhs_pack_nxk_f32p2vlx1biasf32_f32_f32_sme_asm.S
${KLEIDIAI_SRC}/kai/kai_common_sme_asm.S)
set(PRIVATE_ARCH_FLAGS "-fno-tree-vectorize;${PRIVATE_ARCH_FLAGS}+sve+sve2+sme2+fp16")
endif()
+6
View File
@@ -2863,6 +2863,12 @@ struct ggml_cplan ggml_graph_plan(
cur = ggml_type_size(GGML_TYPE_F32) * node->src[0]->ne[0] * n_tasks;
}
} break;
case GGML_OP_SET_ROWS:
{
if (node->src[0]->type == GGML_TYPE_F16 && node->type != GGML_TYPE_F16) {
cur = ggml_type_size(GGML_TYPE_F32) * node->src[0]->ne[0] * n_tasks;
}
} break;
case GGML_OP_SOFT_MAX:
case GGML_OP_ROPE:
case GGML_OP_ROPE_BACK:
+82
View File
@@ -20,14 +20,17 @@
#include "kai_matmul_clamp_f32_qsi8d32p4x8_qsi4c32p8x8_16x8_sve_i8mm.h"
#include "kai_matmul_clamp_f32_qsi8d32p1x8_qsi4c32p8x8_1x8_sve_dotprod.h"
#include "kai_matmul_clamp_f32_f16p1vlx2_qsi4c32p4vlx2_1vlx4vl_sme2_mopa.h"
#include "kai_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa.h"
#include "kai_lhs_pack_bf16p2vlx2_f32_sme.h"
#include "kai_lhs_pack_f32p2vlx1_f32_sme.h"
#include "kai_lhs_quant_pack_qsi8d32p_f32.h"
#include "kai_lhs_quant_pack_qsi8d32p4x8sb_f32_neon.h"
#include "kai_lhs_quant_pack_qsi8d32p_f32_neon.h"
#include "kai_lhs_quant_pack_qai8dxp_f32.h"
#include "kai_rhs_pack_kxn_bf16p2vlx2b_f32_x32_sme.h"
#include "kai_rhs_pack_nxk_f32p2vlx1biasf32_f32_f32_sme.h"
#include "kai_rhs_pack_nxk_qsi4c32pscalef16_qsu4c32s16s0.h"
#include "kai_rhs_pack_nxk_qsi4c32ps1s0scalef16_qsu4c32s16s0_neon.h"
#include "kai_rhs_pack_nxk_qsi8cxp_qsi8cx_neon.h"
@@ -865,6 +868,65 @@ static ggml_kleidiai_kernels gemm_gemv_kernels_q8[] = {
{ /* Sentinel */ }
};
static ggml_kleidiai_kernels ggml_kleidiai_kernels_f32[] = {
#if defined(__ARM_FEATURE_SME)
{
/* SME GEMM */
{
/* .get_m_step = */ kai_get_m_step_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa,
/* .get_n_step = */ kai_get_n_step_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa,
/* .get_mr = */ kai_get_mr_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa,
/* .get_nr = */ kai_get_nr_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa,
/* .get_kr = */ kai_get_kr_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa,
/* .get_sr = */ kai_get_sr_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa,
/* .get_dst_offset = */ kai_get_dst_offset_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa,
/* .get_dst_size = */ kai_get_dst_size_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa,
/* .get_lhs_offset_ex = */ &kernel_offs_fn2<kai_get_lhs_packed_offset_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa>,
/* .get_rhs_packed_offset_ex = */ &kernel_offs_fn2<kai_get_rhs_packed_offset_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa>,
/* .run_kernel_ex = */ &kernel_run_fn10<kai_run_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa>,
},
/* .gemm_lhs_info = */ {
/* .get_offset = */ kai_get_lhs_offset_lhs_pack_f32p2vlx1_f32_sme,
/* .get_packed_offset_ex = */ &lhs_offs_fn5<kai_get_lhs_packed_offset_lhs_pack_f32p2vlx1_f32_sme>,
/* .packed_size_ex = */ &lhs_ps_fn5<kai_get_lhs_packed_size_lhs_pack_f32p2vlx1_f32_sme>,
/* .pack_func_ex = */ &lhs_pack_void_fn9<kai_run_lhs_pack_f32p2vlx1_f32_sme>,
},
/* SME GEMV */
{
/* .get_m_step = */ kai_get_m_step_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa,
/* .get_n_step = */ kai_get_n_step_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa,
/* .get_mr = */ kai_get_mr_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa,
/* .get_nr = */ kai_get_nr_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa,
/* .get_kr = */ kai_get_kr_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa,
/* .get_sr = */ kai_get_sr_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa,
/* .get_dst_offset = */ kai_get_dst_offset_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa,
/* .get_dst_size = */ kai_get_dst_size_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa,
/* .get_lhs_offset_ex = */ nullptr,
/* .get_rhs_packed_offset_ex = */ nullptr,
/* .run_kernel_ex = */ nullptr,
},
/* .gemv_lhs_info = */ {
/* .get_offset = */ kai_get_lhs_offset_lhs_pack_f32p2vlx1_f32_sme,
/* .get_packed_offset_ex = */ &lhs_offs_fn5<kai_get_lhs_packed_offset_lhs_pack_f32p2vlx1_f32_sme>,
/* .packed_size_ex = */ &lhs_ps_fn5<kai_get_lhs_packed_size_lhs_pack_f32p2vlx1_f32_sme>,
/* .pack_func_ex = */ &lhs_pack_void_fn9<kai_run_lhs_pack_f32p2vlx1_f32_sme>,
},
/* .rhs_info = */ {
/* .packed_stride = */ nullptr,
/* .to_float = */ nullptr,
/* .packed_size_ex = */ &rhs_ps_fn2<kai_get_rhs_packed_size_rhs_pack_nxk_f32p2vlx1biasf32_f32_f32_sme>,
/* .packed_stride_ex = */ &rhs_stride_fn1<kai_get_rhs_packed_stride_rhs_pack_nxk_f32p2vlx1biasf32_f32_f32_sme>,
/* .pack_func_ex = */ &rhs_pack_fn13<kai_run_rhs_pack_nxk_f32p2vlx1biasf32_f32_f32_sme>,
},
/* .required_cpu = */ CPU_FEATURE_SME,
/* .lhs_type = */ GGML_TYPE_F32,
/* .rhs_type = */ GGML_TYPE_F32,
/* .op_type = */ GGML_TYPE_F32,
},
#endif
{ /* Sentinel */ }
};
ggml_kleidiai_kernels * ggml_kleidiai_select_kernels(cpu_feature cpu_features, const ggml_tensor * tensor) {
ggml_kleidiai_kernels * kernel = nullptr;
@@ -888,12 +950,15 @@ ggml_kleidiai_kernels * ggml_kleidiai_select_kernels(cpu_feature cpu_features, c
if (tensor->src[0]->type == GGML_TYPE_Q8_0) {
try_table(gemm_gemv_kernels_q8);
} else if (tensor->src[0]->type == GGML_TYPE_F32) {
try_table(ggml_kleidiai_kernels_f32);
} else {
try_table(gemm_gemv_kernels);
}
#else
GGML_UNUSED(gemm_gemv_kernels);
GGML_UNUSED(gemm_gemv_kernels_q8);
GGML_UNUSED(ggml_kleidiai_kernels_f32);
GGML_UNUSED(cpu_features);
#endif
}
@@ -937,3 +1002,20 @@ ggml_kleidiai_kernels * ggml_kleidiai_select_kernels_q8_0(cpu_feature features)
return kernels;
}
ggml_kleidiai_kernels * ggml_kleidiai_select_kernels_f32(cpu_feature features) {
ggml_kleidiai_kernels * kernels = nullptr;
#if defined(__ARM_FEATURE_SME)
for (size_t i = 0; i < NELEMS(ggml_kleidiai_kernels_f32) - 1; ++i) {
if ((features & ggml_kleidiai_kernels_f32[i].required_cpu) == ggml_kleidiai_kernels_f32[i].required_cpu) {
kernels = &ggml_kleidiai_kernels_f32[i];
break;
}
}
#else
GGML_UNUSED(features);
#endif
return kernels;
}
+9
View File
@@ -55,6 +55,12 @@ struct lhs_packing_info {
size_t m_idx_start, const void * lhs, size_t lhs_stride, void * lhs_packed);
};
enum rhs_repack_mode {
RHS_REPACK_PER_KERNEL,
RHS_REPACK_SHARED,
RHS_REPACK_SINGLE_ONLY,
};
struct rhs_packing_info {
size_t (*packed_stride)(size_t k, size_t nr, size_t kr, size_t bl);
@@ -68,6 +74,8 @@ struct rhs_packing_info {
void (*pack_func_ex)(size_t num_groups, size_t n, size_t k, size_t nr, size_t kr, size_t sr, size_t bl,
size_t rhs_stride, const void * rhs, const void * bias, const void * scale, void * rhs_packed, size_t extra_bytes, const void * params);
rhs_repack_mode repack_mode = RHS_REPACK_PER_KERNEL;
};
struct ggml_kleidiai_kernels {
@@ -88,3 +96,4 @@ struct ggml_kleidiai_kernels {
ggml_kleidiai_kernels * ggml_kleidiai_select_kernels(cpu_feature cpu_features, const ggml_tensor * tensor);
ggml_kleidiai_kernels * ggml_kleidiai_select_kernels_q4_0(cpu_feature features);
ggml_kleidiai_kernels * ggml_kleidiai_select_kernels_q8_0(cpu_feature features);
ggml_kleidiai_kernels * ggml_kleidiai_select_kernels_f32(cpu_feature features);
+277 -57
View File
@@ -60,10 +60,11 @@ struct ggml_kleidiai_context {
cpu_feature features;
ggml_kleidiai_kernels * kernels_q4;
ggml_kleidiai_kernels * kernels_q8;
ggml_kleidiai_kernels * kernels_f32;
int sme_thread_cap; // <= 0 means “SME disabled/unknown”;
int thread_hint; // <= 0 means “no hint”
int chunk_multiplier;
} static ctx = { CPU_FEATURE_NONE, nullptr, nullptr, 0, -1, 4 };
} static ctx = { CPU_FEATURE_NONE, nullptr, nullptr, nullptr, 0, -1, 4 };
static const char* cpu_feature_to_string(cpu_feature f) {
if (f == CPU_FEATURE_NONE) {
@@ -156,10 +157,10 @@ static size_t detect_num_smcus() {
}
}
}
return 1;
return 0;
#else
return 1;
return 0;
#endif
}
@@ -192,7 +193,6 @@ static void init_kleidiai_context(void) {
const char *env_threads = getenv("GGML_TOTAL_THREADS");
const char *env_chunk_mult = getenv("GGML_KLEIDIAI_CHUNK_MULTIPLIER");
const bool cpu_has_sme = ggml_cpu_has_sme();
size_t detected_smcus = 0;
ctx.features = (ggml_cpu_has_dotprod() ? CPU_FEATURE_DOTPROD : CPU_FEATURE_NONE) |
@@ -216,56 +216,47 @@ static void init_kleidiai_context(void) {
}
// SME policy:
// - If CPU doesn't support SME: SME always off.
// - Else:
// - env unset => auto-detect cores; enable if detected > 0.
// - env=0 => force off.
// - env>0 => force N cores (skip detection).
// - env unset => auto-detect SMCUs; enable SME only if detected > 0.
// - env=0 => force off.
// - env>0 => force N cores, if the binary was built with SME.
int sme_cores = 0;
bool sme_env_ok = false;
bool sme_env_set = (env_sme != nullptr);
if (!cpu_has_sme) {
if (sme_env_set) {
bool ok = false;
int req = parse_uint_env(env_sme, "GGML_KLEIDIAI_SME", &ok);
if (ok && req > 0) {
GGML_LOG_WARN("kleidiai: GGML_KLEIDIAI_SME=%d but SME is not supported on this CPU; disabling SME\n", req);
}
}
sme_cores = 0;
} else {
if (sme_env_set) {
bool ok = false;
int v = parse_uint_env(env_sme, "GGML_KLEIDIAI_SME", &ok);
sme_env_ok = ok;
if (sme_env_set) {
bool ok = false;
int v = parse_uint_env(env_sme, "GGML_KLEIDIAI_SME", &ok);
sme_env_ok = ok;
if (!ok) {
GGML_LOG_WARN("kleidiai: GGML_KLEIDIAI_SME set but parsing failed; falling back to runtime SME-core detection\n");
detected_smcus = detect_num_smcus();
sme_cores = detected_smcus > 0 ? (int)detected_smcus : 0;
} else if (v == 0) {
sme_cores = 0;
} else {
sme_cores = v;
}
} else {
if (!ok) {
GGML_LOG_WARN("kleidiai: GGML_KLEIDIAI_SME set but parsing failed; falling back to runtime SME-core detection\n");
detected_smcus = detect_num_smcus();
sme_cores = detected_smcus > 0 ? (int)detected_smcus : 0;
} else if (v == 0) {
sme_cores = 0;
} else if (!ggml_cpu_has_sme()) {
GGML_LOG_WARN("kleidiai: GGML_KLEIDIAI_SME=%d but the binary was not built with SME; disabling SME\n", v);
sme_cores = 0;
} else {
sme_cores = v;
}
} else {
detected_smcus = detect_num_smcus();
sme_cores = detected_smcus > 0 ? (int)detected_smcus : 0;
}
if (!sme_env_set && sme_cores == 0) {
GGML_LOG_WARN("kleidiai: SME supported but runtime SME-core detection returned 0; falling back to NEON\n");
}
if (!sme_env_set && ggml_cpu_has_sme() && sme_cores == 0) {
GGML_LOG_WARN("kleidiai: runtime SME-core detection returned 0; falling back to NEON\n");
}
if (sme_cores > 0) {
ctx.features |= CPU_FEATURE_SME;
}
if (sme_cores > 0) {
ctx.features |= CPU_FEATURE_SME;
}
// Kernel selection
ctx.kernels_q4 = ggml_kleidiai_select_kernels_q4_0(ctx.features);
ctx.kernels_q8 = ggml_kleidiai_select_kernels_q8_0(ctx.features);
ctx.kernels_q4 = ggml_kleidiai_select_kernels_q4_0(ctx.features);
ctx.kernels_q8 = ggml_kleidiai_select_kernels_q8_0(ctx.features);
ctx.kernels_f32 = ggml_kleidiai_select_kernels_f32(ctx.features);
if (!ctx.kernels_q4) {
GGML_LOG_INFO("kleidiai: no compatible q4 kernels found for CPU features mask %d\n", (int)ctx.features);
@@ -279,6 +270,12 @@ static void init_kleidiai_context(void) {
GGML_LOG_INFO("kleidiai: primary q8 kernel feature %s\n", cpu_feature_to_string(ctx.kernels_q8->required_cpu));
}
if (!ctx.kernels_f32) {
GGML_LOG_INFO("kleidiai: no compatible f32 kernels found for CPU features mask %d\n", (int)ctx.features);
} else {
GGML_LOG_INFO("kleidiai: primary f32 kernel feature %s\n", cpu_feature_to_string(ctx.kernels_f32->required_cpu));
}
ctx.sme_thread_cap = (ctx.features & CPU_FEATURE_SME) ? sme_cores : 0;
if (ctx.features & CPU_FEATURE_SME) {
@@ -334,6 +331,13 @@ static inline size_t ceil_div_size(size_t a, size_t b) {
return b == 0 ? 0 : (a + b - 1) / b;
}
static inline size_t kleidiai_chunk_cols(size_t n, int nth_total, bool disable_chunking, size_t n_step) {
const size_t multiplier = (nth_total == 1 || disable_chunking) ? 1 : std::max<size_t>(1, (size_t) ctx.chunk_multiplier);
const size_t divisor = std::max<size_t>(1, (size_t) nth_total * multiplier);
const size_t chunk_cols = align_up(std::max<size_t>(1, ceil_div_size(n, divisor)), n_step);
return chunk_cols ? chunk_cols : n_step;
}
struct kleidiai_block_args {
size_t lhs_bl;
size_t rhs_bl;
@@ -418,6 +422,10 @@ static inline ggml_kleidiai_kernels * kleidiai_primary_kernel_q8() {
return ctx.kernels_q8;
}
static inline ggml_kleidiai_kernels * kleidiai_primary_kernel_f32() {
return ctx.kernels_f32;
}
template <typename SelectFallback>
static int kleidiai_collect_kernel_chain_common(
ggml_kleidiai_kernels * primary,
@@ -430,11 +438,16 @@ static int kleidiai_collect_kernel_chain_common(
}
out[count++] = primary;
if (primary->rhs_info.repack_mode == RHS_REPACK_SINGLE_ONLY) {
return count;
}
if ((primary->required_cpu & CPU_FEATURE_SME) == CPU_FEATURE_SME) {
const cpu_feature fallback_mask = static_cast<cpu_feature>(features & ~CPU_FEATURE_SME);
if (fallback_mask != CPU_FEATURE_NONE) {
ggml_kleidiai_kernels * fallback = select_fallback(fallback_mask);
if (fallback && fallback != primary &&
fallback->rhs_info.repack_mode != RHS_REPACK_SINGLE_ONLY &&
fallback->lhs_type == primary->lhs_type &&
fallback->rhs_type == primary->rhs_type &&
fallback->op_type == primary->op_type) {
@@ -465,6 +478,12 @@ static int kleidiai_collect_q8_chain(std::array<ggml_kleidiai_kernels *, GGML_KL
[&](cpu_feature mask) { return ggml_kleidiai_select_kernels_q8_0(mask); });
}
static int kleidiai_collect_f32_chain(std::array<ggml_kleidiai_kernels *, GGML_KLEIDIAI_MAX_KERNEL_SLOTS> & out) {
ggml_kleidiai_kernels * primary = kleidiai_primary_kernel_f32();
return kleidiai_collect_kernel_chain_common(primary, ctx.features, out,
[&](cpu_feature mask) { return ggml_kleidiai_select_kernels_f32(mask); });
}
static inline int64_t ggml_ne(const ggml_tensor * tensor, int dim) {
GGML_ASSERT(dim >= 0 && dim < GGML_MAX_DIMS);
return tensor->ne[dim];
@@ -539,6 +558,36 @@ class tensor_traits : public ggml::cpu::tensor_traits {
return true;
}
if (op->src[0]->type == GGML_TYPE_F32) {
size_t cursor = 0;
bool any_slot = false;
for (int slot = 0; slot < slot_count; ++slot) {
ggml_kleidiai_kernels * kernels = kernel_chain[slot];
lhs_packing_info * lhs_info = &kernels->gemm_lhs_info;
kernel_info * kernel = &kernels->gemm;
if (!lhs_info || !lhs_info->packed_size_ex || !kernel) {
return false;
}
const size_t mr = kernel->get_mr();
const size_t kr = kernel->get_kr();
const size_t sr = kernel->get_sr();
cursor = align_up(cursor, GGML_KLEIDIAI_PACK_ALIGN);
cursor += lhs_info->packed_size_ex(m, k, 0, mr, kr, sr);
any_slot = true;
}
if (!any_slot) {
return false;
}
size = cursor;
return true;
}
if (op->src[0]->type == GGML_TYPE_F16) {
const int64_t lhs_batch_size0 = op->src[1]->ne[2];
const int64_t rhs_batch_size0 = op->src[0]->ne[2];
@@ -595,6 +644,8 @@ class tensor_traits : public ggml::cpu::tensor_traits {
if (dst->op == GGML_OP_MUL_MAT) {
if (dst->src[0]->type == GGML_TYPE_Q4_0 || dst->src[0]->type == GGML_TYPE_Q8_0) {
return compute_forward_qx(params, dst);
} else if (dst->src[0]->type == GGML_TYPE_F32) {
return compute_forward_f32(params, dst);
} else if (dst->src[0]->type == GGML_TYPE_F16) {
return compute_forward_fp16(params, dst);
}
@@ -606,6 +657,144 @@ class tensor_traits : public ggml::cpu::tensor_traits {
return false;
}
bool compute_forward_f32(ggml_compute_params * params, struct ggml_tensor * dst) {
GGML_ASSERT(dst->src[0]->type == GGML_TYPE_F32);
const ggml_tensor * src0 = dst->src[0];
const ggml_tensor * src1 = dst->src[1];
GGML_TENSOR_BINARY_OP_LOCALS
if (src1->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) {
return false;
}
ggml_kleidiai_kernels * kernels = kleidiai_primary_kernel_f32();
if (!kernels) {
return false;
}
kernel_info * kernel = &kernels->gemm;
lhs_packing_info * lhs_info = &kernels->gemm_lhs_info;
if (!kernel || !lhs_info || !lhs_info->get_offset || !lhs_info->get_packed_offset_ex ||
!lhs_info->packed_size_ex || !lhs_info->pack_func_ex ||
!kernel->get_rhs_packed_offset_ex || !kernel->run_kernel_ex || !kernel->get_dst_offset) {
return false;
}
const kleidiai_weight_header * header = kleidiai_weight_header_from_ptr(src0->data);
const bool has_header = kleidiai_is_weight_header_valid(header);
const uint8_t * rhs_base = has_header ? kleidiai_weight_slot_ptr(header, 0)
: static_cast<const uint8_t *>(src0->data);
if (!rhs_base) {
return false;
}
const int nth = params->nth > 0 ? params->nth : 1;
const int ith = params->ith;
const size_t k = ne00;
const size_t m = ne11;
const size_t n = ne01;
const size_t mr = kernel->get_mr();
const size_t kr = kernel->get_kr();
const size_t sr = kernel->get_sr();
const size_t lhs_packed_size = lhs_info->packed_size_ex(m, k, 0, mr, kr, sr);
GGML_ASSERT(lhs_packed_size <= params->wsize);
uint8_t * lhs_packed = static_cast<uint8_t *>(params->wdata);
const size_t dst_stride = dst->nb[1];
const size_t n_step = kernel->get_n_step() ? kernel->get_n_step() : 1;
const bool disable_chunking = ggml_is_numa();
GGML_ASSERT(n <= (size_t) INT_MAX);
for (int64_t batch_idx = 0; batch_idx < ne12; ++batch_idx) {
const uint8_t * lhs_batch_base = static_cast<const uint8_t *>(src1->data) + batch_idx * src1->nb[2];
uint8_t * dst_batch_base = static_cast<uint8_t *>(dst->data) + batch_idx * dst->nb[2];
{
const int64_t m_roundup_mr = kai_roundup((int64_t)m, (int64_t)mr);
int64_t max_threads = mr ? (m_roundup_mr / (int64_t)mr) : nth;
max_threads = std::max<int64_t>(1, max_threads);
const int64_t use_threads = std::min<int64_t>(nth, max_threads);
if (ith < use_threads) {
const int64_t num_m_per_thread0 = round_down((size_t)(m_roundup_mr / use_threads), mr);
const int64_t num_m_per_threadN_1 = (int64_t)m - (use_threads - 1) * num_m_per_thread0;
const int64_t m_start = (int64_t)ith * num_m_per_thread0;
const int64_t m_count = (ith == use_threads - 1) ? num_m_per_threadN_1 : num_m_per_thread0;
const size_t base_packed_off = lhs_info->get_packed_offset_ex(m_start, k, 0, mr, kr, sr);
const size_t next_block_off = lhs_info->get_packed_offset_ex(m_start + mr, k, 0, mr, kr, sr);
const size_t row_stride_bytes = mr ? (next_block_off - base_packed_off) / mr : 0;
int64_t remaining = m_count;
int64_t cur = m_start;
while (remaining > 0) {
const int64_t take = std::min<int64_t>((int64_t)m - cur, remaining);
const size_t src_off = lhs_info->get_offset(cur, src1->nb[1]);
const void * src_ptr = lhs_batch_base + src_off;
const size_t dst_off = base_packed_off + (size_t)(cur - m_start) * row_stride_bytes;
void * dst_ptr = lhs_packed + dst_off;
lhs_info->pack_func_ex(take, k, 0, mr, kr, sr, 0, src_ptr, src1->nb[1], dst_ptr);
cur += take;
remaining -= take;
}
}
}
if (ith == 0) {
ggml_threadpool_chunk_set(params->threadpool, 0);
}
ggml_barrier(params->threadpool);
const size_t chunk_cols = kleidiai_chunk_cols(n, nth, disable_chunking, n_step);
GGML_ASSERT(chunk_cols <= (size_t) INT_MAX);
int current_col = ggml_threadpool_chunk_add(params->threadpool, (int) chunk_cols);
while ((size_t) current_col < n) {
const size_t n_start = (size_t) current_col;
const size_t n_to_process = std::min(chunk_cols, n - n_start);
if (n_to_process > 0) {
const size_t lhs_packed_offset = lhs_info->get_packed_offset_ex(0, k, 0, mr, kr, sr);
const size_t rhs_packed_offset = kernel->get_rhs_packed_offset_ex(n_start, k, 0);
const size_t dst_offset = kernel->get_dst_offset(0, n_start, dst_stride);
const void * lhs_ptr = lhs_packed + lhs_packed_offset;
const void * rhs_ptr = rhs_base + rhs_packed_offset;
float * dst_ptr = reinterpret_cast<float *>(dst_batch_base + dst_offset);
kernel->run_kernel_ex(m, n_to_process, k, 0,
lhs_ptr,
rhs_ptr,
dst_ptr,
dst_stride,
sizeof(float),
-FLT_MAX,
FLT_MAX);
}
current_col = ggml_threadpool_chunk_add(params->threadpool, (int) chunk_cols);
}
if (batch_idx != ne12 - 1) {
ggml_barrier(params->threadpool);
}
}
return true;
}
bool compute_forward_fp16(ggml_compute_params * params, struct ggml_tensor * dst) {
const ggml_tensor * src0 = dst->src[0];
const ggml_tensor * src1 = dst->src[1];
@@ -1214,7 +1403,7 @@ class tensor_traits : public ggml::cpu::tensor_traits {
public:
int repack(struct ggml_tensor * tensor, const void * data, size_t data_size) {
GGML_ASSERT(tensor->type == GGML_TYPE_Q4_0 || tensor->type == GGML_TYPE_Q8_0);
GGML_ASSERT(tensor->type == GGML_TYPE_Q4_0 || tensor->type == GGML_TYPE_Q8_0 || tensor->type == GGML_TYPE_F32);
const size_t n = tensor->ne[1];
const size_t k = tensor->ne[0];
@@ -1233,12 +1422,15 @@ public:
std::array<ggml_kleidiai_kernels *, GGML_KLEIDIAI_MAX_KERNEL_SLOTS> kernel_chain;
const bool want_q8 = tensor->type == GGML_TYPE_Q8_0;
const int slot_total = want_q8 ? kleidiai_collect_q8_chain(kernel_chain)
: kleidiai_collect_q4_chain(kernel_chain);
const bool want_f32 = tensor->type == GGML_TYPE_F32;
const int slot_total = want_f32 ? kleidiai_collect_f32_chain(kernel_chain)
: want_q8 ? kleidiai_collect_q8_chain(kernel_chain)
: kleidiai_collect_q4_chain(kernel_chain);
const bool allow_fallback = kleidiai_pack_fallback_allowed();
std::vector<int8_t> qdata;
std::vector<float> scales;
std::vector<float> bias;
if (want_q8 && slot_total > 0) {
qdata.resize(n * k, 0);
@@ -1286,6 +1478,10 @@ public:
}
}
if (want_f32 && slot_total > 0) {
bias.resize(n, 0.0f);
}
for (int slot = 0; slot < slot_total && slot < GGML_KLEIDIAI_MAX_KERNEL_SLOTS; ++slot) {
if (!allow_fallback && slot > 0) {
break;
@@ -1302,8 +1498,9 @@ public:
const size_t sr = kernel->get_sr();
const ggml_type rhs_type = kernels->rhs_type;
const size_t block_len = rhs_type == GGML_TYPE_Q8_0 ? QK8_0 :
rhs_type == GGML_TYPE_Q4_0 ? QK4_0 : 0;
if (block_len == 0) {
rhs_type == GGML_TYPE_Q4_0 ? QK4_0 :
rhs_type == GGML_TYPE_F32 ? 0 : SIZE_MAX;
if (block_len == SIZE_MAX) {
continue;
}
@@ -1326,6 +1523,10 @@ public:
rhs_info->pack_func_ex(1, n, k, nr, kr, sr, 0, 0,
qdata.data(), nullptr, scales.data(),
dst_ptr, 0, &params);
} else if (rhs_type == GGML_TYPE_F32) {
rhs_info->pack_func_ex(1, n, k, nr, kr, sr, 0, tensor->nb[1],
data, bias.data(), nullptr,
dst_ptr, 0, nullptr);
} else {
continue;
}
@@ -1400,7 +1601,7 @@ static size_t ggml_backend_cpu_kleidiai_buffer_type_get_alignment(ggml_backend_b
static size_t ggml_backend_cpu_kleidiai_buffer_type_get_alloc_size(ggml_backend_buffer_type_t buft, const struct ggml_tensor * tensor) {
GGML_UNUSED(buft);
if (tensor->type != GGML_TYPE_Q4_0 && tensor->type != GGML_TYPE_Q8_0) {
if (tensor->type != GGML_TYPE_Q4_0 && tensor->type != GGML_TYPE_Q8_0 && tensor->type != GGML_TYPE_F32) {
return ggml_nbytes(tensor);
}
@@ -1412,8 +1613,10 @@ static size_t ggml_backend_cpu_kleidiai_buffer_type_get_alloc_size(ggml_backend_
std::array<ggml_kleidiai_kernels *, GGML_KLEIDIAI_MAX_KERNEL_SLOTS> kernel_chain;
const bool want_q8 = tensor->type == GGML_TYPE_Q8_0;
const int slot_total = want_q8 ? kleidiai_collect_q8_chain(kernel_chain)
: kleidiai_collect_q4_chain(kernel_chain);
const bool want_f32 = tensor->type == GGML_TYPE_F32;
const int slot_total = want_f32 ? kleidiai_collect_f32_chain(kernel_chain)
: want_q8 ? kleidiai_collect_q8_chain(kernel_chain)
: kleidiai_collect_q4_chain(kernel_chain);
const bool allow_fallback = kleidiai_pack_fallback_allowed();
size_t slot_count = 0;
@@ -1433,8 +1636,9 @@ static size_t ggml_backend_cpu_kleidiai_buffer_type_get_alloc_size(ggml_backend_
const ggml_type rhs_type = kernels->rhs_type;
const size_t block_len = rhs_type == GGML_TYPE_Q4_0 ? QK4_0 :
rhs_type == GGML_TYPE_Q8_0 ? QK8_0 : 0;
if (block_len == 0) {
rhs_type == GGML_TYPE_Q8_0 ? QK8_0 :
rhs_type == GGML_TYPE_F32 ? 0 : SIZE_MAX;
if (block_len == SIZE_MAX) {
continue;
}
@@ -1455,25 +1659,41 @@ class extra_buffer_type : ggml::cpu::extra_buffer_type {
bool supports_op(ggml_backend_dev_t, const struct ggml_tensor * op) override {
std::array<ggml_kleidiai_kernels *, GGML_KLEIDIAI_MAX_KERNEL_SLOTS> kernel_chain;
const int slot_total = kleidiai_collect_kernel_chain(op, kernel_chain);
if ((op->op == GGML_OP_MUL_MAT || op->op == GGML_OP_GET_ROWS) &&
(op->src[0]->type == GGML_TYPE_Q4_0 || op->src[0]->type == GGML_TYPE_Q8_0) &&
const bool src0_is_kleidiai =
op->src[0]->buffer &&
(ggml_n_dims(op->src[0]) == 2) &&
op->src[0]->buffer->buft == ggml_backend_cpu_kleidiai_buffer_type() &&
slot_total > 0) {
slot_total > 0;
if ((op->op == GGML_OP_MUL_MAT || op->op == GGML_OP_GET_ROWS) &&
(op->src[0]->type == GGML_TYPE_Q4_0 || op->src[0]->type == GGML_TYPE_Q8_0 || op->src[0]->type == GGML_TYPE_F32) &&
src0_is_kleidiai) {
if (op->src[0]->type == GGML_TYPE_Q4_0 && ctx.kernels_q4 == nullptr) {
return false;
}
if (op->src[0]->type == GGML_TYPE_Q8_0 && ctx.kernels_q8 == nullptr) {
return false;
}
if (op->src[0]->type == GGML_TYPE_F32 && ctx.kernels_f32 == nullptr) {
return false;
}
if (op->src[1]->buffer && !ggml_backend_buft_is_host(op->src[1]->buffer->buft)) {
return false;
}
if ((op->src[1]->type == GGML_TYPE_F32 || op->src[1]->type == GGML_TYPE_I32) &&
ggml_ne(op->src[1], 3) == 1) {
return true;
if (op->src[0]->type == GGML_TYPE_Q4_0 || op->src[0]->type == GGML_TYPE_Q8_0) {
if ((op->src[1]->type == GGML_TYPE_F32 || op->src[1]->type == GGML_TYPE_I32) &&
ggml_ne(op->src[1], 3) == 1) {
return true;
}
return false;
}
if (op->op != GGML_OP_MUL_MAT || op->src[1]->type != GGML_TYPE_F32 || op->type != GGML_TYPE_F32) {
return false;
}
return true;
}
return false;
}
+16 -14
View File
@@ -5041,7 +5041,7 @@ static void ggml_compute_forward_set_rows_impl(
assert(ne0 == nc);
assert(ne2 == ne02);
assert(ne3 == ne03);
GGML_ASSERT(src0->type == GGML_TYPE_F32 || (src0->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F16));
GGML_ASSERT(src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16);
assert(ne02 % ne11 == 0);
assert(ne03 % ne12 == 0);
@@ -5075,10 +5075,19 @@ static void ggml_compute_forward_set_rows_impl(
(const float *) ((char *) src0->data + i*nb01 + i02*nb02 + i03*nb03),
((char *) dst->data + i1*nb1 + i02*nb2 + i03*nb3), nc);
} else if constexpr (std::is_same_v<src_t, ggml_fp16_t>) {
memcpy(
if (dst->type == GGML_TYPE_F16) {
memcpy(
((char *) dst->data + i1*nb1 + i02*nb2 + i03*nb3),
((char *) src0->data + i*nb01 + i02*nb02 + i03*nb03),
rs);
} else {
float * wdata = (float *) params->wdata + (nc + CACHE_LINE_SIZE_F32) * ith;
ggml_fp16_to_fp32_row(
(const ggml_fp16_t *) ((char *) src0->data + i*nb01 + i02*nb02 + i03*nb03),
wdata, nc);
from_float(wdata,
((char *) dst->data + i1*nb1 + i02*nb2 + i03*nb3), nc);
}
} else {
GGML_ABORT("src0->type = %d (%s) not supported", src0->type, ggml_type_name(src0->type));
}
@@ -5107,16 +5116,12 @@ void ggml_compute_forward_set_rows(
} break;
case GGML_TYPE_F16:
{
if (dst->type == GGML_TYPE_F16) {
if (src1->type == GGML_TYPE_I64) {
ggml_compute_forward_set_rows_impl<ggml_fp16_t, int64_t>(params, dst);
} else if (src1->type == GGML_TYPE_I32) {
ggml_compute_forward_set_rows_impl<ggml_fp16_t, int32_t>(params, dst);
} else {
GGML_ABORT("src1->type = %d (%s) not supported", src1->type, ggml_type_name(src1->type));
}
if (src1->type == GGML_TYPE_I64) {
ggml_compute_forward_set_rows_impl<ggml_fp16_t, int64_t>(params, dst);
} else if (src1->type == GGML_TYPE_I32) {
ggml_compute_forward_set_rows_impl<ggml_fp16_t, int32_t>(params, dst);
} else {
GGML_ABORT("dst->type = %d (%s) not supported with src0->type = %d (%s)", dst->type, ggml_type_name(dst->type), src0->type, ggml_type_name(src0->type));
GGML_ABORT("src1->type = %d (%s) not supported", src1->type, ggml_type_name(src1->type));
}
} break;
default:
@@ -6362,7 +6367,6 @@ static void ggml_compute_forward_im2col_f16(
const ggml_tensor * src0 = dst->src[0];
const ggml_tensor * src1 = dst->src[1];
GGML_ASSERT(src0->type == GGML_TYPE_F16);
GGML_ASSERT(src1->type == GGML_TYPE_F16 || src1->type == GGML_TYPE_F32);
GGML_ASSERT( dst->type == GGML_TYPE_F16);
@@ -6393,7 +6397,6 @@ static void ggml_compute_forward_im2col_f16(
int ofs0 = is_2D ? nb13 : nb12;
int ofs1 = is_2D ? nb12 : nb11;
GGML_ASSERT(nb00 == sizeof(ggml_fp16_t));
GGML_ASSERT(nb10 == ggml_type_size(src1->type));
// im2col: [N, IC, IH, IW] => [N, OH, OW, IC*KH*KW]
@@ -6563,7 +6566,6 @@ static void ggml_compute_forward_im2col_3d_f16(
const ggml_tensor * src0 = dst->src[0];
const ggml_tensor * src1 = dst->src[1];
GGML_ASSERT(src0->type == GGML_TYPE_F16);
GGML_ASSERT(src1->type == GGML_TYPE_F32);
GGML_ASSERT( dst->type == GGML_TYPE_F16);
+1 -1
View File
@@ -38,7 +38,7 @@ static inline void hmx_queue_process(struct hmx_queue *q, bool* killed) {
if (!d->done) {
FARF(HIGH, "hmx-queue-process: ir %u func %p data %p", ir, d->func, d->data);
enum hmx_queue_signal sig = (enum hmx_queue_signal) (unsigned int) d->func;
uintptr_t sig = (uintptr_t) d->func;
switch (sig) {
case HMX_QUEUE_NOOP: /* noop */; break;
case HMX_QUEUE_KILL: *killed = true; break;
+5 -1
View File
@@ -1340,7 +1340,11 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
return op->src[0]->type != GGML_TYPE_NVFP4;
case GGML_OP_SET_ROWS:
{
if (op->src[0]->type != GGML_TYPE_F32 && op->src[0]->type != GGML_TYPE_F16) {
if (op->src[0]->type == GGML_TYPE_F16) {
return op->type == GGML_TYPE_F16;
}
if (op->src[0]->type != GGML_TYPE_F32) {
return false;
}
+74 -58
View File
@@ -532,6 +532,7 @@ struct ggml_backend_opencl_context {
bool fp16_support;
bool has_vector_subgroup_broadcast;
bool has_subgroup_shuffle = false; // cl_khr_subgroup_shuffle or cl_qcom_subgroup_shuffle
bool has_integer_dot = false; // cl_khr_integer_dot_product or cl_qcom_dot_product8
bool has_qcom_subgroup_shuffle = false; // specifically cl_qcom_subgroup_shuffle
bool disable_fusion;
@@ -834,7 +835,7 @@ struct ggml_backend_opencl_context {
cl_kernel kernel_gemv_moe_q5_1_f32_ns, kernel_gemm_moe_q5_1_f32_ns;
cl_kernel kernel_gemv_moe_q4_k_f32_ns, kernel_gemm_moe_q4_k_f32_ns, kernel_gemm_moe_q4_k_f32_ns_bin;
cl_kernel kernel_gemv_moe_q4_k_f32_ns_wimg = nullptr; // weight-as-texture MoE decode GEMV (opt-in)
cl_kernel kernel_gemm_moe_q4_k_q8_1_dp4a; // dp4a (int8) prefill GEMM variant
cl_kernel kernel_gemm_moe_q4_k_q8_1_dp4a = nullptr; // dp4a (int8) prefill GEMM variant
cl_kernel kernel_moe_reorder_quant_a_q8_1; // fused reorder + q8_1 quant for the dp4a GEMM
cl_kernel kernel_gemm_moe_q8_1_dp4a_q80 = nullptr; // generic dp4a MoE GEMM (MOE_QT=80), opt-in
cl_kernel kernel_moe_expand_scale_q8_0 = nullptr; // q8_0 per-block d -> uniform scale[16]
@@ -844,12 +845,12 @@ struct ggml_backend_opencl_context {
cl_kernel kernel_moe_expand_scale_q5_K = nullptr; // q5_K 6-bit s[] -> uniform scale[16]/min[8]
cl_kernel kernel_gemv_moe_q5_k_f32_ns, kernel_gemm_moe_q5_k_f32_ns;
cl_kernel kernel_gemv_moe_q6_k_f32_ns, kernel_gemm_moe_q6_k_f32_ns;
cl_kernel kernel_gemm_moe_q6_k_q8_1_dp4a; // dp4a (int8) q6_K MoE prefill GEMM
cl_kernel kernel_gemm_moe_q6_k_q8_1_dp4a = nullptr; // dp4a (int8) q6_K MoE prefill GEMM
cl_kernel kernel_gemv_moe_mxfp4_f32, kernel_gemm_moe_mxfp4_f32;
cl_kernel kernel_gemv_moe_mxfp4_f32_ns, kernel_gemm_moe_mxfp4_f32_ns, kernel_gemm_moe_mxfp4_f32_ns_bin;
cl_kernel kernel_gemv_moe_mxfp4_f32_ns_wimg = nullptr; // weight-as-texture MoE decode GEMV
cl_kernel kernel_gemm_moe_mxfp4_q8_1_dp4a; // dp4a (int8) mxfp4 MoE prefill GEMM
cl_kernel kernel_gemm_moe_q4_0_q8_1_dp4a; // dp4a (int8) q4_0 MoE prefill GEMM
cl_kernel kernel_gemm_moe_mxfp4_q8_1_dp4a = nullptr; // dp4a (int8) mxfp4 MoE prefill GEMM
cl_kernel kernel_gemm_moe_q4_0_q8_1_dp4a = nullptr; // dp4a (int8) q4_0 MoE prefill GEMM
cl_kernel kernel_moe_reorder_b;
cl_kernel kernel_moe_histogram, kernel_moe_scan, kernel_moe_fill, kernel_moe_scatter;
cl_kernel kernel_moe_combine_f32 = nullptr; // fused router-weight mul + cross-expert sum
@@ -1037,10 +1038,10 @@ struct ggml_backend_opencl_context {
cl_kernel kernel_gemv_noshuffle_q1_0_f32;
cl_kernel kernel_gemv_noshuffle_q4_k_f32;
cl_kernel kernel_gemm_noshuffle_q4_k_f32;
cl_kernel kernel_gemm_noshuffle_q4_k_q8_1_dp4a; // dp4a (int8) dense prefill GEMM
cl_kernel kernel_gemm_noshuffle_q4_k_q8_1_dp4a_wimg; // dp4a dense prefill GEMM, weights via texture (X1 opt-in)
cl_kernel kernel_gemm_noshuffle_q5_k_q8_1_dp4a; // dp4a (int8) dense q5_K prefill GEMM
cl_kernel kernel_gemm_noshuffle_q6_k_q8_1_dp4a; // dp4a (int8) dense q6_K prefill GEMM
cl_kernel kernel_gemm_noshuffle_q4_k_q8_1_dp4a = nullptr; // dp4a (int8) dense prefill GEMM
cl_kernel kernel_gemm_noshuffle_q4_k_q8_1_dp4a_wimg = nullptr; // dp4a dense prefill GEMM, weights via texture (X1 opt-in)
cl_kernel kernel_gemm_noshuffle_q5_k_q8_1_dp4a = nullptr; // dp4a (int8) dense q5_K prefill GEMM
cl_kernel kernel_gemm_noshuffle_q6_k_q8_1_dp4a = nullptr; // dp4a (int8) dense q6_K prefill GEMM
cl_kernel kernel_quant_a_q8_1; // plain activation q8_1 pre-pass
cl_kernel kernel_gemv_noshuffle_q6_K_f32;
cl_kernel kernel_gemm_noshuffle_q6_K_f32;
@@ -3490,7 +3491,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
}
// gemm_noshuffle_q5_0_q8_1_dp4a (dp4a dense q5_0 prefill GEMM)
{
if (backend_ctx->has_integer_dot) {
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "gemm_noshuffle_q5_0_q8_1_dp4a.cl.h"
@@ -3580,7 +3581,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
}
// gemm_noshuffle_iq4_nl_q8_1_dp4a (dp4a dense IQ4_NL prefill GEMM)
{
if (backend_ctx->has_integer_dot) {
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "gemm_noshuffle_iq4_nl_q8_1_dp4a.cl.h"
@@ -3595,7 +3596,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
}
// gemm_noshuffle_q4_0_q8_1_dp4a (dp4a dense q4_0 prefill GEMM)
{
if (backend_ctx->has_integer_dot) {
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "gemm_noshuffle_q4_0_q8_1_dp4a.cl.h"
@@ -3708,7 +3709,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
}
// gemm_noshuffle_q4_k_q8_1_dp4a (dp4a dense prefill GEMM)
{
if (backend_ctx->has_integer_dot) {
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "gemm_noshuffle_q4_k_q8_1_dp4a.cl.h"
@@ -3730,7 +3731,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
}
// gemm_noshuffle_q8_0_q8_1_dp4a (dp4a dense q8_0 prefill GEMM)
{
if (backend_ctx->has_integer_dot) {
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "gemm_noshuffle_q8_0_q8_1_dp4a.cl.h"
@@ -3746,7 +3747,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
}
// gemm_noshuffle_q5_k_q8_1_dp4a (dp4a dense prefill GEMM for q5_K)
{
if (backend_ctx->has_integer_dot) {
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "gemm_noshuffle_q5_k_q8_1_dp4a.cl.h"
@@ -3761,7 +3762,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
}
// gemm_noshuffle_q6_k_q8_1_dp4a (dp4a dense prefill GEMM for q6_K ffn_down/output)
{
if (backend_ctx->has_integer_dot) {
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "gemm_noshuffle_q6_k_q8_1_dp4a.cl.h"
@@ -4091,7 +4092,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
}
// gemm_moe_q4_k_q8_1_dp4a (dp4a prefill GEMM)
{
if (backend_ctx->has_integer_dot) {
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "gemm_moe_q4_k_q8_1_dp4a.cl.h"
@@ -4108,7 +4109,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
}
// gemm_moe_mxfp4_q8_1_dp4a (dp4a prefill GEMM)
{
if (backend_ctx->has_integer_dot) {
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "gemm_moe_mxfp4_q8_1_dp4a.cl.h"
@@ -4125,7 +4126,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
}
// gemm_moe_q4_0_q8_1_dp4a (dp4a prefill GEMM)
{
if (backend_ctx->has_integer_dot) {
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "gemm_moe_q4_0_q8_1_dp4a.cl.h"
@@ -4142,7 +4143,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
}
// gemm_moe_q8_1_dp4a (generic dp4a MoE GEMM; MOE_QT=80 -> q8_0 expert variant)
{
if (backend_ctx->has_integer_dot) {
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "gemm_moe_q8_1_dp4a.cl.h"
@@ -4256,7 +4257,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
}
// gemm_moe_q6_k_q8_1_dp4a (dp4a q6_K MoE prefill GEMM)
{
if (backend_ctx->has_integer_dot) {
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "gemm_moe_q6_k_q8_1_dp4a.cl.h"
@@ -5602,6 +5603,8 @@ static void ggml_opencl_print_backend_info(ggml_backend_opencl_device_context *
backend_ctx->has_subgroup_shuffle ? "true" : "false");
GGML_LOG_INFO("ggml_opencl: device FP16 support: %s\n",
backend_ctx->fp16_support ? "true" : "false");
GGML_LOG_INFO("ggml_opencl: khr dot product support: %s\n",
backend_ctx->has_integer_dot ? "true" : "false");
GGML_LOG_INFO("ggml_opencl: mem base addr align: %u\n",
backend_ctx->alignment);
GGML_LOG_INFO("ggml_opencl: global mem size: %zu MB\n",
@@ -5810,6 +5813,12 @@ static ggml_backend_opencl_context * ggml_cl_init(ggml_backend_dev_t dev) {
strstr(ext_buffer, "cl_khr_subgroup_shuffle") != NULL ||
backend_ctx->has_qcom_subgroup_shuffle;
// check for cl_khr_integer_dot_product
// cl_qcom_dot_product8 uses signed * unsigned
// while cl_khr_integer_dot_product uses signed * signed -- we stick with khr for now
backend_ctx->has_integer_dot =
strstr(ext_buffer, "cl_khr_integer_dot_product") != NULL;
cl_uint base_align_in_bits;
CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_MEM_BASE_ADDR_ALIGN, sizeof(cl_uint), &base_align_in_bits, NULL));
GGML_ASSERT(base_align_in_bits % 8u == 0);
@@ -10516,10 +10525,16 @@ static ggml_backend_buffer_t ggml_backend_opencl_buffer_type_alloc_buffer(ggml_b
cl_int err;
cl_mem mem = clCreateBuffer(backend_ctx->context, CL_MEM_READ_WRITE, size, NULL, &err);
#if GGML_OPENCL_TARGET_VERSION >= 300
// clCreateBufferWithProperties and cl_mem_properties are OpenCL 3.0. Drivers older than
// that do not export the symbol, so a build targeting them fails to link. The large
// buffer extension is only ever enabled on drivers that are well past 3.0, so this path
// is dead there anyway.
if (err != CL_SUCCESS && backend_ctx->adreno_use_large_buffer) {
cl_mem_properties props[] = { 0x41A6 /* CL_LARGE_BUFFER_QCOM */, 1, 0 };
mem = clCreateBufferWithProperties(backend_ctx->context, props, CL_MEM_READ_WRITE, size, NULL, &err);
}
#endif
if (err != CL_SUCCESS) {
GGML_LOG_INFO("%s: failed to allocate %.2f MiB\n", __func__, size / 1024.0 / 1024.0);
@@ -15819,18 +15834,14 @@ static void ggml_cl_mul_mat_q4_0_f32_adreno(ggml_backend_t backend, const ggml_t
CL_CHECK(clReleaseMemObject(b_sub_buf));
CL_CHECK(clReleaseMemObject(b_img));
} else {
// dp4a (int8) dense prefill GEMM: quant activations to q8_1, then the int8
// dp4a inner-loop GEMM, in place of the transpose + f16 half-dot kernel.
// q4_0 = d*(q-8); mirrors the IQ4_NL/q8_0 dense dp4a paths (+ the sum term).
// OPT-IN / DEFAULT OFF: correct, but neutral on X2E. q4_0's dequant
// ((q-8)*scale) is already trivial so the f16 GEMM is weight-BW-bound and the
// int8 ALU win has nothing to beat -- same as q5_0 dense (unlike IQ4_NL, whose
// codebook dequant is expensive enough for dp4a to help). Kept for A/B; force
// on with GGML_OPENCL_Q4_0_DENSE_DP4A=1. Needs N>8, K%32==0, M%64==0.
// dp4a (int8) dense prefill GEMM, default off
static const char * q4_0_dense_dp4a_env = getenv("GGML_OPENCL_Q4_0_DENSE_DP4A");
const bool q4_0_dense_dp4a_on = q4_0_dense_dp4a_env
bool q4_0_dense_dp4a_on = q4_0_dense_dp4a_env
? (atoi(q4_0_dense_dp4a_env) != 0)
: false;
// dot prod has to be available
q4_0_dense_dp4a_on = backend_ctx->has_integer_dot && q4_0_dense_dp4a_on;
if (q4_0_dense_dp4a_on && backend_ctx->kernel_gemm_noshuffle_q4_0_q8_1_dp4a
&& N > 8 && (K % 32 == 0) && (M % 64 == 0)) {
cl_mem a_sub = nullptr;
@@ -16253,27 +16264,16 @@ static void ggml_cl_mul_mat_q5_0_f32_adreno(ggml_backend_t backend, const ggml_t
CL_CHECK(clReleaseMemObject(b_sub_buf));
CL_CHECK(clReleaseMemObject(b_img));
} else {
// dp4a (int8) dense q5_0 prefill GEMM. Quantizes the [N,K] activations to
// q8_1 and runs the int8 dot instead of the f16 half-dot. Large-batch
// (ne1>8) only. q5_0 weight = (x-16)*d (x = nibble | hi<<4); x packed as a
// 0..31 byte (dp4a), the -16 centering folded into a single min term
// (d*16) via the q8_1 block sum. Reads the qs/qh/d buffers byte-identically
// to the f16 kernel (greedy byte-identical, MUL_MAT NMSE-OK).
//
// OPT-IN / DEFAULT OFF. Unlike q8_0/q4_K dense, dp4a is not a win for q5_0 on
// X2E: the q5_0 model is bottlenecked elsewhere, so the dense-GEMM int8 win
// has nothing to surface and the q8_1 prepass slightly hurts. Kept correct +
// opt-in for the X1 A/B (different texture-cache dynamic) and the
// weight-texture variant. Env: GGML_OPENCL_Q5_DENSE_DP4A=1.
// Weight-as-texture variant (X1 lever): routes the dominant qs nibble plane
// through an image1d_buffer (qh stays a buffer). Opt-in
// GGML_OPENCL_Q5_DENSE_DP4A_WIMG; when set it also forces the dp4a path on.
// dp4a (int8) dense q5_0 prefill GEMM, default off
static const char * q5_dense_dp4a_env = getenv("GGML_OPENCL_Q5_DENSE_DP4A");
static const char * q5_dense_wimg_env = getenv("GGML_OPENCL_Q5_DENSE_DP4A_WIMG");
const bool q5_dense_wimg_on = q5_dense_wimg_env && (atoi(q5_dense_wimg_env) != 0);
const bool q5_dense_dp4a_on = q5_dense_wimg_on
bool q5_dense_dp4a_on = q5_dense_wimg_on
? true
: (q5_dense_dp4a_env && (atoi(q5_dense_dp4a_env) != 0));
// dot prod has to be available
q5_dense_dp4a_on = backend_ctx->has_integer_dot && q5_dense_dp4a_on;
if (q5_dense_dp4a_on && backend_ctx->kernel_gemm_noshuffle_q5_0_q8_1_dp4a
&& N > 8 && (K % 32 == 0) && (M % 64 == 0)) {
cl_mem a_sub = nullptr;
@@ -16708,15 +16708,14 @@ static void ggml_cl_mul_mat_iq4_nl_f32_adreno(ggml_backend_t backend, const ggml
} else {
// dp4a (int8) dense IQ4_NL prefill GEMM. Quantizes the [N,K] activations to
// q8_1 and runs the int8 dot instead of the f16 half-dot. Large-batch
// (ne1>8) only. IQ4_NL weight = kvalues[nibble]*d; the codebook value IS the
// int8 (no min term), so this is the q8_0 dense case plus a nibble->int8 LUT
// unpack. Reads the q/d buffers byte-identically to the f16 kernel. No bin
// kernel for IQ4_NL -> baseline is f16, default ON for X2E (like q4_K/q6_K
// dense dp4a). X1 stays on f16. Env: GGML_OPENCL_IQ4NL_DENSE_DP4A.
// (ne1>8) only
static const char * iq4nl_dense_dp4a_env = getenv("GGML_OPENCL_IQ4NL_DENSE_DP4A");
const bool iq4nl_dense_dp4a_on = iq4nl_dense_dp4a_env
bool iq4nl_dense_dp4a_on = iq4nl_dense_dp4a_env
? (atoi(iq4nl_dense_dp4a_env) != 0)
: (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E);
// dot prod has to be available
iq4nl_dense_dp4a_on = backend_ctx->has_integer_dot && iq4nl_dense_dp4a_on;
if (iq4nl_dense_dp4a_on && backend_ctx->kernel_gemm_noshuffle_iq4_nl_q8_1_dp4a
&& N > 8 && (K % 32 == 0) && (M % 64 == 0)) {
cl_mem a_sub = nullptr;
@@ -16966,13 +16965,15 @@ static void ggml_cl_mul_mat_q8_0_f32_adreno(ggml_backend_t backend, const ggml_t
static const char * q8_dense_wimg_env = getenv("GGML_OPENCL_Q8_DENSE_DP4A_WIMG");
const bool q8_dense_wimg_on = q8_dense_wimg_env && (atoi(q8_dense_wimg_env) != 0);
const bool q8_bin_loaded = (backend_ctx->kernel_gemm_noshuffle_q8_0_f32_bin != nullptr);
const bool q8_bin_loaded = (backend_ctx->kernel_gemm_noshuffle_q8_0_f32_bin != nullptr);
// bin kernel takes precedence
const bool q8_dense_dp4a_on = q8_dense_wimg_on
bool q8_dense_dp4a_on = q8_dense_wimg_on
? true
: q8_dense_dp4a_env
? (atoi(q8_dense_dp4a_env) != 0)
: (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E && !q8_bin_loaded);
// dot prod has to be available
q8_dense_dp4a_on = backend_ctx->has_integer_dot && q8_dense_dp4a_on;
if (q8_dense_dp4a_on && backend_ctx->kernel_gemm_noshuffle_q8_0_q8_1_dp4a
&& N > 8 && (K % 32 == 0) && (M % 64 == 0)) {
@@ -17379,13 +17380,16 @@ static void ggml_cl_mul_mat_q4_k_f32_adreno(ggml_backend_t backend, const ggml_t
static const char * q4k_dense_dp4a_env = getenv("GGML_OPENCL_Q4K_DENSE_DP4A");
static const char * q4k_dense_wimg_env = getenv("GGML_OPENCL_Q4K_DENSE_DP4A_WIMG");
const bool q4k_dense_wimg_on = q4k_dense_wimg_env && (atoi(q4k_dense_wimg_env) != 0);
const bool q4k_dense_dp4a_on = q4k_dense_wimg_on
const bool q4k_dense_wimg_on = q4k_dense_wimg_env && (atoi(q4k_dense_wimg_env) != 0);
bool q4k_dense_dp4a_on = q4k_dense_wimg_on
? true
: q4k_dense_dp4a_env
? (atoi(q4k_dense_dp4a_env) != 0)
: (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E);
// dp4 has to be available
q4k_dense_dp4a_on = backend_ctx->has_integer_dot && q4k_dense_dp4a_on;
// Min N for the dp4a prefill GEMM, default 9, i.e., ne1 > 8
static const char * q4k_dp4a_minn_env = getenv("GGML_OPENCL_Q4K_DP4A_MINN");
const int q4k_dp4a_minn = q4k_dp4a_minn_env ? atoi(q4k_dp4a_minn_env) : 9;
@@ -17608,9 +17612,11 @@ static void ggml_cl_mul_mat_q6_K_f32_adreno(ggml_backend_t backend, const ggml_t
// dp4a (int8) dense q6_K prefill GEMM
static const char * q6k_dense_dp4a_env = getenv("GGML_OPENCL_Q6K_DENSE_DP4A");
static const bool q6k_dense_dp4a_on = (q6k_dense_dp4a_env != nullptr)
bool q6k_dense_dp4a_on = (q6k_dense_dp4a_env != nullptr)
? (atoi(q6k_dense_dp4a_env) != 0)
: (backend_ctx->adreno_gen != ADRENO_GPU_GEN::X1E);
// dot prod has to be available
q6k_dense_dp4a_on = backend_ctx->has_integer_dot && q6k_dense_dp4a_on;
const bool is_output_w_dp4a = strncmp(src0->name, "output", 6) == 0 ||
strncmp(src0->name, "token_embd", 10) == 0;
@@ -17901,9 +17907,11 @@ static void ggml_cl_mul_mat_q5_K_f32_adreno(ggml_backend_t backend, const ggml_t
// dp4a (int8) dense q5_K prefill GEMM
static const char * q5k_dense_dp4a_env = getenv("GGML_OPENCL_Q5K_DENSE_DP4A");
const bool q5k_dense_dp4a_on = q5k_dense_dp4a_env
bool q5k_dense_dp4a_on = q5k_dense_dp4a_env
? (atoi(q5k_dense_dp4a_env) != 0)
: (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E);
// dot prod has to be available
q5k_dense_dp4a_on = backend_ctx->has_integer_dot && q5k_dense_dp4a_on;
if (q5k_dense_dp4a_on && ne1 > 8 && (ne00 % 32 == 0) && (ne01 % 64 == 0)) {
const int Mm = ne01, Nn = ne1, Kk = ne00;
@@ -20640,6 +20648,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
bool use_moe_dp4a = q4_0_moe_dp4a_env
? (atoi(q4_0_moe_dp4a_env) != 0)
: (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E);
// dot prod has to be available
use_moe_dp4a = backend_ctx->has_integer_dot && use_moe_dp4a;
// bin kernel takes precedence
use_moe_dp4a = use_moe_dp4a && backend_ctx->kernel_gemm_moe_q4_0_f32_ns_bin == nullptr;
@@ -21815,6 +21825,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
bool use_moe_dp4a = (q4k_moe_dp4a_env != nullptr)
? (atoi(q4k_moe_dp4a_env) != 0)
: (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E || backend_ctx->adreno_gen == ADRENO_GPU_GEN::X1E);
// dot prod has to be available
use_moe_dp4a = backend_ctx->has_integer_dot && use_moe_dp4a;
// bin kernel takes precedence
use_moe_dp4a = use_moe_dp4a && backend_ctx->kernel_gemm_moe_q4_k_f32_ns_bin == nullptr;
@@ -22316,10 +22328,12 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
// dp4a (int8) q6_K MoE prefill GEMM
static const char * q6k_moe_dp4a_env = getenv("GGML_OPENCL_Q6K_MOE_DP4A");
static const bool use_moe_dp4a = (q6k_moe_dp4a_env != nullptr)
bool use_moe_dp4a = (q6k_moe_dp4a_env != nullptr)
? (atoi(q6k_moe_dp4a_env) != 0)
: (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E
|| backend_ctx->adreno_gen == ADRENO_GPU_GEN::X1E);
// dot prod has to be available
use_moe_dp4a = backend_ctx->has_integer_dot && use_moe_dp4a;
cl_buffer_region region;
region.origin = 0;
@@ -22569,6 +22583,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
bool use_moe_dp4a = mxfp4_moe_dp4a_env
? (atoi(mxfp4_moe_dp4a_env) != 0)
: (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E);
// dot prod has to be available
use_moe_dp4a = backend_ctx->has_integer_dot && use_moe_dp4a;
// bin kernel takes precedence
use_moe_dp4a = use_moe_dp4a && backend_ctx->kernel_gemm_moe_mxfp4_f32_ns_bin == nullptr;
@@ -296,7 +296,12 @@ kernel void kernel_gemv_noshuffle_iq4_nl_f32(
// 2 outputs per fiber in wave 0
if (groupId == 0) {
dst = (global float*)((global char*)dst + offsetd);
vstore2(totalSum, 0, &(dst[gid * 2]));
// Guard the two output rows. The x-grid is padded to CEIL_DIV(ne01/2,64)*64,
// so when ne01 is not a multiple of 128 the tail row-pairs run past row ne01
// and would overrun dst into the adjacent tensor. No-op / byte-identical when
// ne01 % 128 == 0 (M/2 already a multiple of 64 -> no padding).
if (gid * 2 + 0 < M) dst[gid * 2 + 0] = totalSum.s0;
if (gid * 2 + 1 < M) dst[gid * 2 + 1] = totalSum.s1;
}
}
@@ -116,6 +116,10 @@ __kernel void kernel_gemv_noshuffle_q1_0_f32(
if (groupId == 0) {
dst = (global float*)((global char*)dst + offsetd);
dst[gid] = totalSum;
// Guard the output row. The x-grid is padded to CEIL_DIV(M,wavesize)*wavesize,
// so when ne01 is not a multiple of the wave size the tail work-items run past
// row ne01 and would overrun dst into the adjacent tensor. No-op / byte-identical
// when ne01 is wave-aligned (no padding).
if (gid < M) dst[gid] = totalSum;
}
}
@@ -268,7 +268,12 @@ __kernel void kernel_gemv_noshuffle_q4_0_f32(
// 2 outputs per fiber in wave 0
if (groupId == 0) {
dst = (global float*)((global char*)dst + offsetd);
vstore2(totalSum, 0, &(dst[gid * 2]));
// Guard the two output rows. The x-grid is padded to CEIL_DIV(ne01/2,64)*64,
// so when ne01 is not a multiple of 128 the tail row-pairs run past row ne01
// and would overrun dst into the adjacent tensor. No-op / byte-identical when
// ne01 % 128 == 0 (M/2 already a multiple of 64 -> no padding).
if (gid * 2 + 0 < M) dst[gid * 2 + 0] = totalSum.s0;
if (gid * 2 + 1 < M) dst[gid * 2 + 1] = totalSum.s1;
}
}
@@ -262,7 +262,11 @@ __kernel void kernel_gemv_noshuffle_q4_0_f32(
// 2 outputs per fiber in wave 0
if (groupId == 0) {
dst = (global float*)((global char*)dst + offsetd);
vstore2(totalSum, 0, &(dst[gid * 2]));
// Guard the two output rows against the padded x-grid tail overrunning dst.
// The current shape specializations are all ne01 % 128 == 0 (no padding), so
// this is a no-op / byte-identical today; keep it in lockstep with the base kernel.
if (gid * 2 + 0 < ne01) dst[gid * 2 + 0] = totalSum.s0;
if (gid * 2 + 1 < ne01) dst[gid * 2 + 1] = totalSum.s1;
}
}
@@ -277,7 +277,12 @@ kernel void kernel_gemv_noshuffle_q4_1_f32(
// 2 outputs per fiber in wave 0
if (groupId == 0) {
dst = (global float*)((global char*)dst + offsetd);
vstore2(totalSum, 0, &(dst[gid * 2]));
// Guard the two output rows. The x-grid is padded to CEIL_DIV(ne01/2,64)*64,
// so when ne01 is not a multiple of 128 the tail row-pairs run past row ne01
// and would overrun dst into the adjacent tensor. No-op / byte-identical when
// ne01 % 128 == 0 (M/2 already a multiple of 64 -> no padding).
if (gid * 2 + 0 < M) dst[gid * 2 + 0] = totalSum.s0;
if (gid * 2 + 1 < M) dst[gid * 2 + 1] = totalSum.s1;
}
}
@@ -312,7 +312,12 @@ kernel void kernel_gemv_noshuffle_q4_k_f32(
// 2 outputs per fiber in wave 0
if (groupId == 0) {
dst = (global float*)((global char*)dst + offsetd);
vstore2(totalSum, 0, &(dst[gid * 2]));
// Guard the two output rows. The x-grid is padded to CEIL_DIV(ne01/2,64)*64,
// so when ne01 is not a multiple of 128 the tail row-pairs run past row ne01
// and would overrun dst into the adjacent tensor. No-op / byte-identical when
// ne01 % 128 == 0 (M/2 already a multiple of 64 -> no padding).
if (gid * 2 + 0 < M) dst[gid * 2 + 0] = totalSum.s0;
if (gid * 2 + 1 < M) dst[gid * 2 + 1] = totalSum.s1;
}
}
@@ -285,7 +285,12 @@ __kernel void kernel_gemv_noshuffle_q5_0_f32(
// 2 outputs per fiber in wave 0
if (groupId == 0) {
dst = (global float*)((global char*)dst + offsetd);
vstore2(totalSum, 0, &(dst[gid * 2]));
// Guard the two output rows. The x-grid is padded to CEIL_DIV(ne01/2,64)*64,
// so when ne01 is not a multiple of 128 the tail row-pairs run past row ne01
// and would overrun dst into the adjacent tensor. No-op / byte-identical when
// ne01 % 128 == 0 (M/2 already a multiple of 64 -> no padding).
if (gid * 2 + 0 < M) dst[gid * 2 + 0] = totalSum.s0;
if (gid * 2 + 1 < M) dst[gid * 2 + 1] = totalSum.s1;
}
}
@@ -288,7 +288,12 @@ __kernel void kernel_gemv_noshuffle_q5_1_f32(
// 2 outputs per fiber in wave 0
if (groupId == 0) {
dst = (global float*)((global char*)dst + offsetd);
vstore2(totalSum, 0, &(dst[gid * 2]));
// Guard the two output rows. The x-grid is padded to CEIL_DIV(ne01/2,64)*64,
// so when ne01 is not a multiple of 128 the tail row-pairs run past row ne01
// and would overrun dst into the adjacent tensor. No-op / byte-identical when
// ne01 % 128 == 0 (M/2 already a multiple of 64 -> no padding).
if (gid * 2 + 0 < M) dst[gid * 2 + 0] = totalSum.s0;
if (gid * 2 + 1 < M) dst[gid * 2 + 1] = totalSum.s1;
}
}
@@ -321,6 +321,11 @@ kernel void kernel_gemv_noshuffle_q5_k_f32(
// 2 outputs per fiber in wave 0
if (groupId == 0) {
dst = (global float*)((global char*)dst + offsetd);
vstore2(totalSum, 0, &(dst[gid * 2]));
// Guard the two output rows. The x-grid is padded to CEIL_DIV(ne01/2,64)*64,
// so when ne01 is not a multiple of 128 the tail row-pairs run past row ne01
// and would overrun dst into the adjacent tensor. No-op / byte-identical when
// ne01 % 128 == 0 (M/2 already a multiple of 64 -> no padding).
if (gid * 2 + 0 < M) dst[gid * 2 + 0] = totalSum.s0;
if (gid * 2 + 1 < M) dst[gid * 2 + 1] = totalSum.s1;
}
}
@@ -288,6 +288,11 @@ kernel void kernel_gemv_noshuffle_q6_K_f32(
if (grp == 0) {
dst = (global float*)((global char*)dst + offsetd);
vstore2(total_sum, 0, &(dst[gid * 2]));
// Guard the two output rows. The x-grid is padded to CEIL_DIV(ne01/2,64)*64,
// so when ne01 is not a multiple of 128 the tail row-pairs run past row ne01
// and would overrun dst into the adjacent tensor (garbage downstream).
// No-op / byte-identical when ne01 % 128 == 0 (no padding).
if (gid * 2 + 0 < ne01) dst[gid * 2 + 0] = total_sum.s0;
if (gid * 2 + 1 < ne01) dst[gid * 2 + 1] = total_sum.s1;
}
}
@@ -190,6 +190,10 @@ __kernel void kernel_gemv_noshuffle_q8_0_f32(
// 1 outputs per fiber in wave 0
if (groupId == 0) {
dst = (global float*)((global char*)dst + offsetd);
dst[gid] = totalSum;
// Guard the output row. The x-grid is padded to CEIL_DIV(M,wavesize)*wavesize,
// so when ne01 is not a multiple of the wave size the tail work-items run past
// row ne01 and would overrun dst into the adjacent tensor. No-op / byte-identical
// when ne01 is wave-aligned (no padding).
if (gid < M) dst[gid] = totalSum;
}
}
@@ -64,7 +64,14 @@ kernel void kernel_mul_mat_f16_f16(
global half * x = (global half *) (src0 + offset_src0);
if (ne00 < 128) {
// The vector path below casts the row pointers to half4, which must be 8-byte aligned.
// A row address is r0*nb01 + ..., and a permuted or strided src leaves nb01/nb11
// unconstrained -- an odd ne00, say, gives a row that is only 2-byte aligned. Every
// src1 row this work-item walks is src1_base + r1*nb11, so require both.
const ulong src1_base = (ulong) (src1 + (i12)*nb12 + (i13)*nb13);
const bool row_aligned = (((ulong) x) & 7) == 0 && (src1_base & 7) == 0 && (nb11 & 7) == 0;
if (ne00 < 128 || !row_aligned) {
for (int row = 0; row < N_F16_F16; ++row) {
int r1 = rb + row;
if (r1 >= ne11) {
@@ -64,7 +64,14 @@ kernel void kernel_mul_mat_f16_f32(
global half * x = (global half *) (src0 + offset_src0);
if (ne00 < 128) {
// The vector path below casts the row pointers to half4/float4, which must be 8- and
// 16-byte aligned. A row address is r0*nb01 + ..., and a permuted or strided src leaves
// nb01/nb11 unconstrained -- an odd ne00, say, gives a row that is only 2-byte aligned.
// Every src1 row this work-item walks is src1_base + r1*nb11, so require both.
const ulong src1_base = (ulong) (src1 + (i12)*nb12 + (i13)*nb13);
const bool row_aligned = (((ulong) x) & 7) == 0 && (src1_base & 15) == 0 && (nb11 & 15) == 0;
if (ne00 < 128 || !row_aligned) {
for (int row = 0; row < N_F16_F32; ++row) {
int r1 = rb + row;
if (r1 >= ne11) {
@@ -64,8 +64,15 @@ kernel void kernel_mul_mat_f16_f32_1row(
global half * x = (global half *) (src0 + offset_src0);
global float * y = (global float *) (src1 + offset_src1);
// The vector path below casts the row pointers to half4/float4, which must be 8- and
// 16-byte aligned. A row address is r0*nb01 + ..., and a permuted or strided src leaves
// nb01/nb11 unconstrained -- an odd ne00, say, gives a row that is only 2-byte aligned.
// Take the vector path only when the rows this work-item touches are actually aligned;
// the scalar loop has no such requirement.
const bool row_aligned = (((ulong) x) & 7) == 0 && (((ulong) y) & 15) == 0;
float sumf = 0;
if (ne00 < 128) {
if (ne00 < 128 || !row_aligned) {
for (int i = get_sub_group_local_id(); i < ne00; i += get_max_sub_group_size()) {
sumf += (float) x[i] * (float) y[i];
}
+1
View File
@@ -64,6 +64,7 @@ extern int g_ggml_sycl_enable_fusion;
extern int g_ggml_sycl_prioritize_dmmv;
extern int g_ggml_sycl_enable_flash_attention;
extern int g_ggml_sycl_dev2dev_memcpy;
extern int g_ggml_sycl_fa_onednn;
#if defined(__clang__) && __has_builtin(__builtin_expect)
+25 -13
View File
@@ -71,8 +71,8 @@ struct dw_cwhn_layout {
}
};
template <typename Layout>
static void conv2d_dw_kernel(const float * input, const float * kernel, float * output,
template <typename KernelT, typename Layout>
static void conv2d_dw_kernel(const float * input, const KernelT * kernel, float * output,
const conv2d_dw_params p, const sycl::nd_item<3> & item_ct1) {
const int global_idx = item_ct1.get_local_id(2) +
item_ct1.get_group(2) * item_ct1.get_local_range(2);
@@ -93,15 +93,15 @@ static void conv2d_dw_kernel(const float * input, const float * kernel, float *
for (int kx = bounds.x_min; kx < bounds.x_max; ++kx) {
const int in_x = dw_calculate_input_coord(out_x, kx, p.stride_x, p.dilation_x, p.padding_x);
acc += input[Layout::input_index(n, c, in_y, in_x, p)] *
kernel[Layout::kernel_index(c, ky, kx, p)];
static_cast<float>(kernel[Layout::kernel_index(c, ky, kx, p)]);
}
}
output[Layout::output_index(n, c, out_y, out_x, p)] = acc;
}
template <typename Layout>
static void conv2d_dw_sycl(const float * x_d, const float * w_d, float * y_d,
template <typename KernelT, typename Layout>
static void conv2d_dw_sycl(const float * x_d, const KernelT * w_d, float * y_d,
const conv2d_dw_params p, const queue_ptr & stream) {
const int total = p.batches * p.channels * p.out_h * p.out_w;
const int num_blocks = (total + SYCL_CONV2D_DW_BLOCK_SIZE - 1) / SYCL_CONV2D_DW_BLOCK_SIZE;
@@ -109,7 +109,7 @@ static void conv2d_dw_sycl(const float * x_d, const float * w_d, float * y_d,
const sycl::range<3> block_nums(1, 1, num_blocks);
stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> item_ct1) {
conv2d_dw_kernel<Layout>(x_d, w_d, y_d, p, item_ct1);
conv2d_dw_kernel<KernelT, Layout>(x_d, w_d, y_d, p, item_ct1);
});
}
@@ -119,9 +119,9 @@ void ggml_sycl_op_conv2d_dw(ggml_backend_sycl_context & ctx, ggml_tensor * dst)
const ggml_tensor * kernel = dst->src[0];
const ggml_tensor * input = dst->src[1];
GGML_ASSERT(kernel->type == GGML_TYPE_F32 && input->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32);
GGML_ASSERT((kernel->type == GGML_TYPE_F32 || kernel->type == GGML_TYPE_F16) &&
input->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32);
const float * w_d = (const float *) kernel->data;
const float * x_d = (const float *) input->data;
float * y_d = (float *) dst->data;
@@ -148,11 +148,23 @@ void ggml_sycl_op_conv2d_dw(ggml_backend_sycl_context & ctx, ggml_tensor * dst)
const queue_ptr stream = ctx.stream();
if (ggml_is_contiguous(input)) {
conv2d_dw_sycl<dw_whcn_layout>(x_d, w_d, y_d, params, stream);
} else if (ggml_is_contiguous_channels(input)) {
conv2d_dw_sycl<dw_cwhn_layout>(x_d, w_d, y_d, params, stream);
if (kernel->type == GGML_TYPE_F16) {
const sycl::half * w_d = (const sycl::half *) kernel->data;
if (ggml_is_contiguous(input)) {
conv2d_dw_sycl<sycl::half, dw_whcn_layout>(x_d, w_d, y_d, params, stream);
} else if (ggml_is_contiguous_channels(input)) {
conv2d_dw_sycl<sycl::half, dw_cwhn_layout>(x_d, w_d, y_d, params, stream);
} else {
GGML_ABORT("Unsupported memory layout for conv2d_dw");
}
} else {
GGML_ABORT("Unsupported memory layout for conv2d_dw");
const float * w_d = (const float *) kernel->data;
if (ggml_is_contiguous(input)) {
conv2d_dw_sycl<float, dw_whcn_layout>(x_d, w_d, y_d, params, stream);
} else if (ggml_is_contiguous_channels(input)) {
conv2d_dw_sycl<float, dw_cwhn_layout>(x_d, w_d, y_d, params, stream);
} else {
GGML_ABORT("Unsupported memory layout for conv2d_dw");
}
}
}
+115 -16
View File
@@ -19,6 +19,7 @@
typedef void (*dequantize_kernel_t)(const void * vx, const int64_t ib, const int iqs, dfloat2 & v);
typedef void (*dequantize_kernel_t_reorder)(const void *d, const int64_t ib, const void *qs,
const int iqs, dfloat2 &v);
typedef void (*dequantize_kernel_f32_t)(const void * vx, const int64_t ib, const int iqs, float & v0, float & v1);
#if QK_K == 256
static inline void get_scale_min_k4(int j, const uint8_t * q, uint8_t & d, uint8_t & m);
@@ -85,6 +86,21 @@ static __dpct_inline__ void dequantize_q1_0_reorder(const void *d_ptr, const int
v.y() = (2 * bit_1 - 1) * d;
}
static __dpct_inline__ void dequantize_q1_0(const void *vx, const int64_t ib,
const int iqs, dfloat2 &v) {
const block_q1_0 * x = (const block_q1_0 *) vx;
const dfloat d = x[ib].d;
const int bit_index_0 = iqs + 0;
const int bit_index_1 = iqs + 1;
const int bit_0 = (x[ib].qs[bit_index_0 / 8] >> (bit_index_0 % 8)) & 1;
const int bit_1 = (x[ib].qs[bit_index_1 / 8] >> (bit_index_1 % 8)) & 1;
v.x() = (2 * bit_0 - 1) * d;
v.y() = (2 * bit_1 - 1) * d;
}
static __dpct_inline__ void dequantize_q4_1(const void *vx, const int64_t ib,
const int iqs, dfloat2 &v) {
const block_q4_1 * x = (const block_q4_1 *) vx;
@@ -140,6 +156,39 @@ static __dpct_inline__ void dequantize_q4_K(const void *vx, const int64_t ib,
#endif
}
static __dpct_inline__ void dequantize_q4_K_f32(const void *vx, const int64_t ib,
const int iqs, float &v0, float &v1) {
#if QK_K == 256
const block_q4_K * x = (const block_q4_K *) vx;
const sycl::half2 dm = x[ib].dm;
const float dall = dm[0];
const float dmin = dm[1];
auto dequantize_one = [&](const int idx) -> float {
const int il = idx / 64;
const int in = idx % 64;
const int is = 2 * il + (in >= 32 ? 1 : 0);
const int qsi = 32 * il + (in & 31);
uint8_t sc;
uint8_t m;
get_scale_min_k4(is, x[ib].scales, sc, m);
const float d = dall * sc;
const float mn = dmin * m;
const uint8_t q = x[ib].qs[qsi];
const uint8_t qv = (in >= 32) ? (q >> 4) : (q & 0xF);
return d * qv - mn;
};
v0 = dequantize_one(iqs + 0);
v1 = dequantize_one(iqs + 1);
#else
GGML_ABORT("Q4_K dequantize not supported for QK_K != 256");
#endif
}
static __dpct_inline__ void dequantize_q2_K(const void *vx, const int64_t ib,
const int iqs, dfloat2 &v) {
#if QK_K == 256
@@ -159,7 +208,7 @@ static __dpct_inline__ void dequantize_q2_K(const void *vx, const int64_t ib,
const float d = dall * (sc & 0xF);
const float m = dmin * (sc >> 4);
return sycl::fma((dfloat) ((q >> (2 * g)) & 3), (dfloat) d, (dfloat) (-m));
return (dfloat) d * (dfloat) ((q >> (2 * g)) & 3) - (dfloat) m;
};
v.x() = dequantize_one(iqs + 0);
@@ -169,6 +218,35 @@ static __dpct_inline__ void dequantize_q2_K(const void *vx, const int64_t ib,
#endif
}
static __dpct_inline__ void dequantize_q2_K_f32(const void *vx, const int64_t ib,
const int iqs, float &v0, float &v1) {
#if QK_K == 256
const block_q2_K * x = (const block_q2_K *) vx;
const float dall = x[ib].dm[0];
const float dmin = x[ib].dm[1];
auto dequantize_one = [&](const int idx) -> float {
const int n = idx / 128;
const int r = idx % 128;
const int g = r / 32;
const int l = r % 32;
const int is = 8 * n + l / 16;
const uint8_t q = x[ib].qs[32 * n + l];
const uint8_t sc = x[ib].scales[is + 2 * g];
const float d = dall * (sc & 0xF);
const float m = dmin * (sc >> 4);
return d * ((q >> (2 * g)) & 3) - m;
};
v0 = dequantize_one(iqs + 0);
v1 = dequantize_one(iqs + 1);
#else
GGML_ABORT("Q2_K dequantize not supported for QK_K != 256");
#endif
}
static __dpct_inline__ void dequantize_q3_K(const void *vx, const int64_t ib,
const int iqs, dfloat2 &v) {
#if QK_K == 256
@@ -242,6 +320,42 @@ static __dpct_inline__ void dequantize_q5_K(const void *vx, const int64_t ib,
#endif
}
static __dpct_inline__ void dequantize_q5_K_f32(const void *vx, const int64_t ib,
const int iqs, float &v0, float &v1) {
#if QK_K == 256
const block_q5_K * x = (const block_q5_K *) vx;
const float dall = x[ib].dm[0];
const float dmin = x[ib].dm[1];
auto dequantize_one = [&](const int idx) -> float {
const int il = idx / 64;
const int in = idx % 64;
const int is = 2 * il + (in >= 32 ? 1 : 0);
const int ir = (in & 31) / 2;
const int iq = in & 1;
const uint8_t q = x[ib].qs[32 * il + 2 * ir + iq];
const uint8_t h = x[ib].qh[2 * ir + iq];
const uint8_t qv = (in >= 32) ? (q >> 4) : (q & 0xF);
uint8_t sc;
uint8_t m;
get_scale_min_k4(is, x[ib].scales, sc, m);
const float d = dall * sc;
const float mn = dmin * m;
const uint8_t hm = 1 << (2 * il + (in >= 32 ? 1 : 0));
return (qv + ((h & hm) ? 16 : 0)) * d - mn;
};
v0 = dequantize_one(iqs + 0);
v1 = dequantize_one(iqs + 1);
#else
GGML_ABORT("Q5_K dequantize not supported for QK_K != 256");
#endif
}
static __dpct_inline__ void dequantize_q6_K(const void *vx, const int64_t ib,
const int iqs, dfloat2 &v) {
#if QK_K == 256
@@ -296,21 +410,6 @@ static __dpct_inline__ void dequantize_mxfp4(const void *vx, const int64_t ib,
v.y() = d * kvalues_mxfp4[q >> 4] * 0.5f;
}
static __dpct_inline__ void dequantize_q1_0(const void *vx, const int64_t ib,
const int iqs, dfloat2 &v) {
const block_q1_0 * x = (const block_q1_0 *) vx;
const dfloat d = x[ib].d;
const int bit_index_0 = iqs + 0;
const int bit_index_1 = iqs + 1;
const int bit_0 = (x[ib].qs[bit_index_0 / 8] >> (bit_index_0 % 8)) & 1;
const int bit_1 = (x[ib].qs[bit_index_1 / 8] >> (bit_index_1 % 8)) & 1;
v.x() = (2 * bit_0 - 1) * d;
v.y() = (2 * bit_1 - 1) * d;
}
static __dpct_inline__ void dequantize_nvfp4(const void *vx, const int64_t ib,
const int iqs, dfloat2 &v) {
const block_nvfp4 & xb = ((const block_nvfp4 *) vx)[ib];
+40
View File
@@ -247,6 +247,17 @@ static __dpct_inline__ T op_leaky_relu(T x, float negative_slope) {
}
}
template<typename T>
static __dpct_inline__ T op_xielu(T x, float alpha_n, float alpha_p, float beta, float eps) {
const float xi = static_cast<float>(x);
const float gate_pos = (xi > 0.0f);
const float y_pos = alpha_p * xi * xi + beta * xi;
const float min_v_eps = sycl::fmin(xi, eps);
const float y_neg = (sycl::expm1(min_v_eps) - xi) * alpha_n + beta * xi;
const float out = gate_pos * y_pos + (1.0f - gate_pos) * y_neg;
return static_cast<T>(out);
}
template<typename T>
static __dpct_inline__ T op_sqr(T x) {
return x * x;
@@ -359,6 +370,13 @@ static void unary_op_leaky_relu_kernel(const T * x, T * dst, const int k, float
}
}
template<typename T>
static void unary_op_xielu_kernel(const T * x, T * dst, const int k, float alpha_n, float alpha_p, float beta, float eps, const sycl::nd_item<1> &item_ct1) {
SYCL_GLOBAL_ID_LOOP(k, item_ct1) {
dst[i] = op_xielu(x[i], alpha_n, alpha_p, beta, eps);
}
}
template<typename T>
static void unary_op_sqr_kernel(const T * x, T * dst, const int k, const sycl::nd_item<1> &item_ct1) {
SYCL_GLOBAL_ID_LOOP(k, item_ct1) {
@@ -836,6 +854,23 @@ static inline void ggml_sycl_op_clamp(ggml_backend_sycl_context & ctx, ggml_tens
}, min_val, max_val);
}
static inline void ggml_sycl_op_xielu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
const float alpha_n = ggml_get_op_params_f32(dst, 1);
const float alpha_p = ggml_get_op_params_f32(dst, 2);
const float beta = ggml_get_op_params_f32(dst, 3);
const float eps = ggml_get_op_params_f32(dst, 4);
ggml_sycl_detail::dispatch_ggml_sycl_op_unary(ctx, dst,
[](const auto* src, auto* dst_ptr, int k_elements, queue_ptr stream, float alpha_n_arg, float alpha_p_arg, float beta_arg, float eps_arg) {
const int num_blocks = ceil_div(k_elements, SYCL_RELU_BLOCK_SIZE);
stream->parallel_for(
sycl::nd_range<1>(sycl::range<1>(num_blocks) * sycl::range<1>(SYCL_RELU_BLOCK_SIZE),
sycl::range<1>(SYCL_RELU_BLOCK_SIZE)),
[=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
unary_op_xielu_kernel(src, dst_ptr, k_elements, alpha_n_arg, alpha_p_arg, beta_arg, eps_arg, item_ct1);
});
}, alpha_n, alpha_p, beta, eps);
}
static inline void ggml_sycl_op_floor(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
ggml_sycl_detail::ggml_sycl_op_unary(ctx, dst, [](auto x) {
return op_floor(x);
@@ -1153,6 +1188,11 @@ void ggml_sycl_clamp(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
ggml_sycl_op_clamp(ctx, dst);
}
void ggml_sycl_xielu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
ggml_sycl_op_xielu(ctx, dst);
}
void ggml_sycl_sgn(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
ggml_sycl_op_sgn(ctx, dst);
+2
View File
@@ -75,6 +75,8 @@ void ggml_sycl_sqr(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
void ggml_sycl_clamp(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
void ggml_sycl_xielu(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
void ggml_sycl_sgn(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
void ggml_sycl_abs(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
+265
View File
@@ -0,0 +1,265 @@
#include <cstdint>
#include <cstdio>
#include <cstring>
#include <string>
#include <unordered_map>
#include <vector>
#include "fattn-onednn.hpp"
#include "fattn-tile.hpp"
// set minimum query length to treat as prefill (32)
#define GGML_SYCL_FA_ONEDNN_MIN_Q 32
bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst) {
#if !GGML_SYCL_DNNL
GGML_UNUSED(dst);
return false;
#else
if (!g_ggml_sycl_fa_onednn) {
return false;
}
// Battlemage (Xe2) only, for now. On other Intel archs oneDNN's fused SDPA returns wrong results
// for some shapes (e.g. head_dim=64 on Arc / xe_hpg) -- an oneDNN bug tracked upstream at
// https://github.com/uxlfoundation/oneDNN/issues/5510. Remove this hardware limitation once that
// is fixed; until then non-BMG archs fall back to the existing FA kernel.
const gpu_arch arch = ggml_sycl_info().devices[ggml_sycl_get_device()].hw_info.arch;
if (arch != gpu_arch::intel_gpu_bmg_g21 && arch != gpu_arch::intel_gpu_bmg_g31) {
return false;
}
const ggml_tensor * Q = dst->src[0];
const ggml_tensor * K = dst->src[1];
const ggml_tensor * V = dst->src[2];
const ggml_tensor * mask = dst->src[3];
const ggml_tensor * sinks = dst->src[4];
// gate for f16 KV only for now
// need to implement quantized KV
if (K->type != GGML_TYPE_F16 || V->type != GGML_TYPE_F16) {
return false;
}
// gate for the following cases
// 1. if the oneDNN graph Add node has no input --> skip
// 2. types other than f16 need different logical_tensor declaration
// 3. the mask must be shape [1, 1, q, seq]
// 4. sinks: excludes attention sink (Xiao et al., 2024) that can't be modeled by oneDNN graph
if (!mask || mask->type != GGML_TYPE_F16 || mask->ne[2] != 1 || mask->ne[3] != 1 || sinks) {
return false;
}
float max_bias = 0.0f, logit_softcap = 0.0f;
memcpy(&max_bias, (const float *) dst->op_params + 1, sizeof(float));
memcpy(&logit_softcap, (const float *) dst->op_params + 2, sizeof(float));
if (max_bias != 0.0f || logit_softcap != 0.0f) {
return false;
}
// K and V must share head_dim: the SDPA graph uses a single `d` for both.
const int64_t d = K->ne[0];
if (V->ne[0] != d || Q->ne[3] != 1) {
return false;
}
// GQA must divide evenly.
if (K->ne[2] == 0 || Q->ne[2] % K->ne[2] != 0) {
return false;
}
// Prefill only.
if (Q->ne[1] < GGML_SYCL_FA_ONEDNN_MIN_Q) {
return false;
}
return true;
#endif
}
#if GGML_SYCL_DNNL
#include "dnnl.hpp"
#include "dnnl_sycl.hpp"
#include "oneapi/dnnl/dnnl_graph.hpp" // graph API lives only under oneapi/dnnl/, not at the include root
using namespace dnnl;
using namespace dnnl::graph;
// strided src (f16 or f32) -> contiguous f16 [ne0,ne1,ne2,ne3] (ne0 innermost). nb* are BYTE strides.
template <typename src_t>
static void cont_to_f16_sycl(const char * src, sycl::half * dst,
int64_t ne0, int64_t ne1, int64_t ne2, int64_t ne3,
size_t nb1, size_t nb2, size_t nb3, dpct::queue_ptr stream) {
const int64_t n = ne0 * ne1 * ne2 * ne3;
stream->parallel_for(sycl::range<1>(n), [=](sycl::id<1> ix) {
const int64_t gid = ix[0];
int64_t i = gid;
const int64_t i0 = i % ne0; i /= ne0;
const int64_t i1 = i % ne1; i /= ne1;
const int64_t i2 = i % ne2; const int64_t i3 = i / ne2;
const src_t * p = (const src_t *) (src + i1 * nb1 + i2 * nb2 + i3 * nb3) + i0;
dst[gid] = (sycl::half) (*p);
});
}
// oneDNN SDPA out (f16 contiguous [mb,H,q,d]) -> ggml dst (f32 [head_dim,H,n_tok,mb], contiguous).
static void permute_sdpa_out_sycl(const sycl::half * out, float * dst,
int64_t mb, int64_t H, int64_t q, int64_t d, dpct::queue_ptr stream) {
const int64_t n = mb * H * q * d;
stream->parallel_for(sycl::range<1>(n), [=](sycl::id<1> ix) {
const int64_t gid = ix[0];
int64_t i = gid;
const int64_t e = i % d; i /= d;
const int64_t t = i % q; i /= q;
const int64_t h = i % H; const int64_t b = i / H;
dst[e + h * d + t * d * H + b * d * H * q] = (float) out[gid];
});
}
struct sdpa_partition {
compiled_partition cp;
std::vector<logical_tensor> ins;
logical_tensor out;
size_t id_q = 0, id_k = 0, id_v = 0, id_scale = 0, id_mask = 0;
bool ok = false;
};
// Build + compile the contiguous-input GQA SDPA graph (MatMul->Divide->Add->SoftMax->MatMul), f32 out.
// Mirrors the hardware-verified scratch/onednn_sdpa_probe.cpp build_gqa (partitions=1, sdp_primitive_kernel_t).
static sdpa_partition build_sdpa(const engine & eng, int H, int Hkv, int q, int seq, int d) {
using ltype = logical_tensor::layout_type;
using dt = logical_tensor::data_type;
using ldims = logical_tensor::dims;
const dt fi = dt::f32, t = dt::f16;
const int rep = H / Hkv;
const ldims q_sz = {1, Hkv, rep, q, d}, kv_sz = {1, Hkv, 1, seq, d}, s_sz = {1, Hkv, rep, q, seq},
sc = {1, 1, 1, 1, 1}, msk = {1, 1, 1, q, seq}, o_sz = {1, Hkv, rep, q, d};
int64_t id = 0;
sdpa_partition E;
auto query = logical_tensor(id++, t, q_sz, ltype::strided);
auto key = logical_tensor(id++, t, kv_sz, ltype::strided);
auto score = logical_tensor(id++, fi, s_sz, ltype::strided);
auto bmm1 = op(id++, op::kind::MatMul, "bmm1");
bmm1.set_attr<bool>(op::attr::transpose_b, true); // key is [.., seq, d]
bmm1.add_inputs({query, key}); bmm1.add_outputs({score});
auto scale = logical_tensor(id++, t, sc, ltype::strided);
auto scaled = logical_tensor(id++, fi, s_sz, ltype::strided);
auto sdiv = op(id++, op::kind::Divide, "scale_div"); // score / (1/kq_scale) == score * kq_scale
sdiv.add_inputs({score, scale}); sdiv.add_outputs({scaled});
auto mask = logical_tensor(id++, t, msk, ltype::strided);
auto masked = logical_tensor(id++, fi, s_sz, ltype::strided);
auto madd = op(id++, op::kind::Add, "mask_add");
madd.add_inputs({scaled, mask}); madd.add_outputs({masked});
auto probs = logical_tensor(id++, t, s_sz, ltype::strided);
auto smax = op(id++, op::kind::SoftMax, "softmax");
smax.set_attr<int64_t>(op::attr::axis, -1);
smax.set_attr<std::string>(op::attr::mode, "inf_as_zero");
smax.add_inputs({masked}); smax.add_outputs({probs});
auto value = logical_tensor(id++, t, kv_sz, ltype::strided);
// f16 output is REQUIRED to hit sdp_primitive_kernel_t (the systolic micro-kernel); an f32 output
// falls to larger_partition_kernel_t which materializes N^2 (confirmed: scratch/onednn_sdpa_kernel_probe.cpp).
// converted to the f32 ggml dst in the permute below.
auto output = logical_tensor(id++, t, o_sz, ltype::strided); // f16 contiguous [mb,Hkv,rep,q,d]
auto bmm2 = op(id++, op::kind::MatMul, "bmm2");
bmm2.add_inputs({probs, value}); bmm2.add_outputs({output});
dnnl::graph::graph g(eng.get_kind());
g.add_op(bmm1); g.add_op(sdiv); g.add_op(madd); g.add_op(smax); g.add_op(bmm2);
g.finalize();
auto parts = g.get_partitions();
if (parts.size() != 1 || !parts[0].is_supported()) {
return E; // ok stays false -> caller falls back to TILE
}
E.ins = parts[0].get_input_ports();
E.out = parts[0].get_output_ports()[0];
E.cp = parts[0].compile(E.ins, {E.out}, eng);
E.out = E.cp.query_logical_tensor(E.out.get_id());
E.id_q = query.get_id(); E.id_k = key.get_id(); E.id_v = value.get_id();
E.id_scale = scale.get_id(); E.id_mask = mask.get_id();
E.ok = true;
return E;
}
void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tensor * dst) try {
const ggml_tensor * Q = dst->src[0];
const ggml_tensor * K = dst->src[1];
const ggml_tensor * V = dst->src[2];
const ggml_tensor * mask = dst->src[3];
const int64_t d = K->ne[0]; // head_dim
const int64_t seq = K->ne[1]; // n_kv
const int64_t Hkv = K->ne[2]; // n_head_kv
const int64_t H = Q->ne[2]; // n_head
const int64_t q = Q->ne[1]; // n_tok
const int64_t mb = Q->ne[3]; // batch (== 1, gated)
float kq_scale = 1.0f;
memcpy(&kq_scale, (const float *) dst->op_params + 0, sizeof(float));
dpct::queue_ptr stream = ctx.stream();
dnnl::engine eng = ctx.engine_dnnl(stream);
dnnl::stream strm = ctx.stream_dnnl(stream);
// cont/cast inputs to contiguous f16 (head-major) -- the layout the fast systolic path wants.
ggml_sycl_pool_alloc<sycl::half> Qf(ctx.pool(), (size_t) H * q * d);
ggml_sycl_pool_alloc<sycl::half> Kf(ctx.pool(), (size_t) Hkv * seq * d);
ggml_sycl_pool_alloc<sycl::half> Vf(ctx.pool(), (size_t) Hkv * seq * d);
cont_to_f16_sycl<float> ((const char *) Q->data, Qf.get(), d, q, H, mb, Q->nb[1], Q->nb[2], Q->nb[3], stream);
cont_to_f16_sycl<sycl::half>((const char *) K->data, Kf.get(), d, seq, Hkv, mb, K->nb[1], K->nb[2], K->nb[3], stream);
cont_to_f16_sycl<sycl::half>((const char *) V->data, Vf.get(), d, seq, Hkv, mb, V->nb[1], V->nb[2], V->nb[3], stream);
// divide-by-(1/scale) reproduces ggml's score *= kq_scale on the proven probe graph.
const sycl::half scale_h = (sycl::half) (1.0f / kq_scale);
ggml_sycl_pool_alloc<sycl::half> scbuf(ctx.pool(), 1);
stream->memcpy(scbuf.get(), &scale_h, sizeof(sycl::half));
ggml_sycl_pool_alloc<sycl::half> outf(ctx.pool(), (size_t) H * q * d); // f16 contiguous SDPA out [mb,H,q,d]
// compile once per (device, shape), reuse across layers/calls.
static std::unordered_map<std::string, sdpa_partition> cache;
char keyb[96];
snprintf(keyb, sizeof(keyb), "%d:%lld:%lld:%lld:%lld:%lld", ggml_sycl_get_device(),
(long long) H, (long long) Hkv, (long long) q, (long long) seq, (long long) d);
auto it = cache.find(keyb);
if (it == cache.end()) {
it = cache.emplace(keyb, build_sdpa(eng, (int) H, (int) Hkv, (int) q, (int) seq, (int) d)).first;
}
sdpa_partition & E = it->second;
// _supported() is authoritative: if it accepted this op the partition must build.
// A failure here is a gap in _supported() -- surface it, don't mask it with a fallback.
GGML_ASSERT(E.ok && "oneDNN SDPA partition failed to build for a _supported() shape");
auto id2ptr = [&](size_t r) -> void * {
if (r == E.id_q) return Qf.get();
if (r == E.id_k) return Kf.get();
if (r == E.id_v) return Vf.get();
if (r == E.id_scale) return scbuf.get();
if (r == E.id_mask) return (void *) mask->data;
return nullptr;
};
std::vector<tensor> ti;
ti.reserve(E.ins.size());
for (auto & lt : E.ins) {
ti.emplace_back(lt, eng, id2ptr(lt.get_id()));
}
tensor to(E.out, eng, outf.get());
E.cp.execute(strm, ti, {to});
permute_sdpa_out_sycl(outf.get(), (float *) dst->data, mb, H, q, d, stream);
// Single device: no sync is required, and actually PP perf is ~6% > wait_and_throw() (tested on llama-3.1-8b & qwen3.6-27b, both Q8_0, with Arc B70).
// Any future multi-GPU refactor MUST re-measure this single-device path and keep the best
// single-device PP speed. Otherwise (multiple devices/streams can race the reuse):
if (ggml_sycl_info().device_count > 1) {
// cont_to_f16 -> oneDNN execute -> permute is async on this stream, but the
// pool_alloc*s above free their device buffers at host return. Without this wait the next
// scheduler op re-acquires those bytes while the GPU is still computing the SDPA, turning
// it into garbage and collapsing multi-turn output to a single repeated token ("GGGGG...").
stream->wait_and_throw();
}
}
catch (const std::exception & e) {
// any oneDNN/SYCL failure is non-fatal: fall back to the existing kernel (strictly additive).
GGML_LOG_WARN("%s: oneDNN SDPA failed (%s); falling back to TILE kernel\n", __func__, e.what());
ggml_sycl_flash_attn_ext_tile(ctx, dst);
}
#endif // GGML_SYCL_DNNL
+14
View File
@@ -0,0 +1,14 @@
#ifndef GGML_SYCL_FATTN_ONEDNN_HPP
#define GGML_SYCL_FATTN_ONEDNN_HPP
#include "common.hpp"
// Static-only check: fused-XMX oneDNN Graph SDPA path==flash-attn op
// (f16 KV, no softcap/ALiBi, single stream, tuned head_dim, prefill-sized q.)
bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst);
// Run flash attention through oneDNN's fused xmx SDPA
// execute the cached SDPA partition, write the f32 dst. Falls back to the TILE kernel on any failure.
void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
#endif // GGML_SYCL_FATTN_ONEDNN_HPP
+14 -1
View File
@@ -18,6 +18,7 @@
#include "fattn-tile.hpp"
#include "fattn-vec.hpp"
#include "fattn.hpp"
#include "fattn-onednn.hpp"
#define FATTN_VEC_CASE(D, type_K, type_V) \
@@ -96,6 +97,7 @@ static void ggml_sycl_flash_attn_ext_vec(ggml_backend_sycl_context & ctx, ggml_t
enum best_fattn_kernel {
BEST_FATTN_KERNEL_NONE = 0,
BEST_FATTN_KERNEL_VEC = 100,
BEST_FATTN_KERNEL_ONEDNN = 150, // added enum for onednn==150
BEST_FATTN_KERNEL_TILE = 200,
};
@@ -189,7 +191,11 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const
// For small batch sizes the vector kernel may be preferable over the kernels optimized for large batch sizes:
const bool can_use_vector_kernel = Q->ne[0] <= 512 && Q->ne[0] % 64 == 0 && K->ne[1] % FATTN_KQ_STRIDE == 0;
// Todo: Use the XMX kernel if possible:
// Fused-XMX path: oneDNN Graph SDPA (flash attention). Strictly
// additive -- taken only when statically supported, otherwise falls through to VEC/TILE below.
if (ggml_sycl_flash_attn_ext_onednn_supported(dst)) {
return BEST_FATTN_KERNEL_ONEDNN;
}
// If there are no tensor cores available, use the generic tile kernel:
if (can_use_vector_kernel) {
@@ -213,6 +219,13 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst
switch (ggml_sycl_get_best_fattn_kernel(ggml_sycl_get_device(), dst)) {
case BEST_FATTN_KERNEL_NONE:
GGML_ABORT("Not support Flash-Attention");
case BEST_FATTN_KERNEL_ONEDNN:
// guarded: ggml_sycl_flash_attn_ext_onednn() is only defined under GGML_SYCL_DNNL;
// the reference must be compiled out here or the GGML_SYCL_DNNL=0 build fails to link.
#if GGML_SYCL_DNNL
ggml_sycl_flash_attn_ext_onednn(ctx, dst);
#endif
break;
case BEST_FATTN_KERNEL_TILE:
ggml_sycl_flash_attn_ext_tile(ctx, dst);
break;
+80 -3
View File
@@ -60,6 +60,50 @@ static void k_get_rows(
dst_row[iybs + iqs + y_offset] = v.y();
}
template<int qk, int qr, dequantize_kernel_f32_t dequantize_kernel, typename dst_t>
static void k_get_rows_f32(
const void * src0, const int32_t * src1, dst_t * dst,
int64_t ne00,
int64_t ne12,
size_t s1, size_t s2, size_t s3,
size_t nb01, size_t nb02, size_t nb03,
size_t s10, size_t s11, size_t s12,
const sycl::nd_item<3> &item_ct1) {
const int i00 = (item_ct1.get_group(2) * item_ct1.get_local_range(2) +
item_ct1.get_local_id(2)) *
2;
const int i10 = item_ct1.get_local_range(1) * item_ct1.get_group(1) +
item_ct1.get_local_id(1);
const int i11 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) +
item_ct1.get_local_id(0)) /
ne12;
const int i12 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) +
item_ct1.get_local_id(0)) %
ne12;
if (i00 >= ne00) {
return;
}
const int i01 = src1[i10*s10 + i11*s11 + i12*s12];
dst_t * dst_row = dst + i10*s1 + i11*s2 + i12*s3;
const void * src0_row = (const char *)src0 + i01*nb01 + i11*nb02 + i12*nb03;
const int ib = i00/qk;
const int iqs = (i00%qk)/qr;
const int iybs = i00 - i00%qk;
const int y_offset = qr == 1 ? 1 : qk/2;
float v0;
float v1;
dequantize_kernel(src0_row, ib, iqs, v0, v1);
dst_row[iybs + iqs + 0] = (dst_t) v0;
dst_row[iybs + iqs + y_offset] = (dst_t) v1;
}
template<typename src0_t, typename dst_t>
static void k_get_rows_float(
const src0_t * src0, const int32_t * src1, dst_t * dst,
@@ -129,6 +173,39 @@ static void get_rows_sycl(ggml_backend_sycl_context & ctx, const ggml_tensor *sr
GGML_UNUSED(ctx);
}
template <int qk, int qr, dequantize_kernel_f32_t dq>
static void get_rows_sycl_f32(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1,
ggml_tensor *dst, const void *src0_dd,
const int32_t *src1_dd, float *dst_dd,
queue_ptr stream) {
GGML_TENSOR_BINARY_OP_LOCALS
const sycl::range<3> block_dims(1, 1, SYCL_GET_ROWS_BLOCK_SIZE);
const int block_num_x = (ne00 + 2*SYCL_GET_ROWS_BLOCK_SIZE - 1) / (2*SYCL_GET_ROWS_BLOCK_SIZE);
const sycl::range<3> block_nums(ne11 * ne12, ne10, block_num_x);
const size_t s1 = nb1 / ggml_element_size(dst);
const size_t s2 = nb2 / ggml_element_size(dst);
const size_t s3 = nb3 / ggml_element_size(dst);
const size_t s10 = nb10 / ggml_element_size(src1);
const size_t s11 = nb11 / ggml_element_size(src1);
const size_t s12 = nb12 / ggml_element_size(src1);
GGML_ASSERT(ne00 % 2 == 0);
stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> item_ct1) {
k_get_rows_f32<qk, qr, dq>(
src0_dd, src1_dd, dst_dd, ne00, ne12, s1, s2,
s3, nb01, nb02, nb03, s10, s11, s12, item_ct1);
});
GGML_UNUSED(dst);
GGML_UNUSED(ctx);
}
template <typename src0_t, typename dst_t>
static void get_rows_sycl_float(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
const ggml_tensor *src1, ggml_tensor *dst,
@@ -244,7 +321,7 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
src1_i32, (float *)dst->data, ctx.stream());
break;
case GGML_TYPE_Q2_K:
get_rows_sycl<QK_K, 1, dequantize_q2_K>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
get_rows_sycl_f32<QK_K, 1, dequantize_q2_K_f32>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
src1_i32, (float *)dst->data, ctx.stream());
break;
case GGML_TYPE_Q3_K:
@@ -260,7 +337,7 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
src1_i32, (float *)dst->data, ctx.stream());
break;
case GGML_TYPE_Q4_K:
get_rows_sycl<QK_K, 1, dequantize_q4_K>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
get_rows_sycl_f32<QK_K, 1, dequantize_q4_K_f32>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
src1_i32, (float *)dst->data, ctx.stream());
break;
case GGML_TYPE_Q5_0:
@@ -272,7 +349,7 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
src1_i32, (float *)dst->data, ctx.stream());
break;
case GGML_TYPE_Q5_K:
get_rows_sycl<QK_K, 1, dequantize_q5_K>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
get_rows_sycl_f32<QK_K, 1, dequantize_q5_K_f32>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
src1_i32, (float *)dst->data, ctx.stream());
break;
case GGML_TYPE_Q6_K:
+10 -1
View File
@@ -84,6 +84,7 @@ int g_ggml_sycl_debug = 0;
int g_ggml_sycl_enable_optimize = 1;
int g_ggml_sycl_enable_graph = 0;
int g_ggml_sycl_enable_dnn = 1;
int g_ggml_sycl_fa_onednn = 1;
int g_ggml_sycl_enable_vmm = 1;
int g_ggml_sycl_enable_fusion = 1;
int g_ggml_sycl_prioritize_dmmv = 0;
@@ -285,6 +286,7 @@ static void ggml_check_sycl() try {
g_ggml_sycl_enable_optimize = ggml_sycl_get_env("GGML_SYCL_ENABLE_OPT", 1);
g_ggml_sycl_enable_graph = ggml_sycl_get_env("GGML_SYCL_ENABLE_GRAPH", 0);
g_ggml_sycl_enable_dnn = ggml_sycl_get_env("GGML_SYCL_ENABLE_DNN", 1);
g_ggml_sycl_fa_onednn = ggml_sycl_get_env("GGML_SYCL_FA_ONEDNN", 1);
g_ggml_sycl_enable_vmm = ggml_sycl_get_env("GGML_SYCL_ENABLE_VMM", 1);
g_ggml_sycl_enable_fusion = ggml_sycl_get_env("GGML_SYCL_ENABLE_FUSION", 1);
g_ggml_sycl_prioritize_dmmv = ggml_sycl_get_env("GGML_SYCL_PRIORITIZE_DMMV", 0);
@@ -352,8 +354,10 @@ static void ggml_check_sycl() try {
#if defined(GGML_SYCL_DNNL)
GGML_LOG_INFO(" GGML_SYCL_ENABLE_DNN: %d\n", g_ggml_sycl_enable_dnn);
GGML_LOG_INFO(" GGML_SYCL_FA_ONEDNN: %d\n", g_ggml_sycl_fa_onednn);
#else
GGML_LOG_INFO(" GGML_SYCL_ENABLE_DNN: DNN disabled by compile flag\n");
GGML_LOG_INFO(" GGML_SYCL_FA_ONEDNN: %d\n", g_ggml_sycl_fa_onednn);
#endif
#ifdef SYCL_FLASH_ATTN
GGML_LOG_INFO(" GGML_SYCL_ENABLE_FLASH_ATTN: %d\n", g_ggml_sycl_enable_flash_attention);
@@ -839,7 +843,7 @@ static const char * ggml_backend_sycl_buffer_type_get_name(ggml_backend_buffer_t
}
static bool check_usm_system(int device, size_t size) {
bool use_usm_system = g_ggml_sycl_usm_system && size >= MEM_SIZE_1G;
bool use_usm_system = g_ggml_sycl_usm_system && size >= ((size_t)4 * MEM_SIZE_1G);
if (use_usm_system && !ggml_sycl_info().devices[device].usm_system_support) {
GGML_LOG_INFO("Device does not support USM system allocations\n");
@@ -878,6 +882,7 @@ ggml_backend_sycl_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft,
void * dev_ptr;
if (use_usm_system) {
GGML_SYCL_DEBUG("[SYCL] allocating %lu Bytes with USM system\n", size);
dev_ptr = (void *)aligned_malloc_host(alignment, aligned_size);
if (!dev_ptr) {
GGML_LOG_ERROR("%s: can't allocate %lu Bytes of memory on host\n", __func__, size);
@@ -5006,6 +5011,9 @@ static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct gg
case GGML_UNARY_OP_ELU:
ggml_sycl_elu(ctx, dst);
break;
case GGML_UNARY_OP_XIELU:
ggml_sycl_xielu(ctx, dst);
break;
case GGML_UNARY_OP_FLOOR:
ggml_sycl_floor(ctx, dst);
break;
@@ -5668,6 +5676,7 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
case GGML_UNARY_OP_EXPM1:
case GGML_UNARY_OP_SOFTPLUS:
case GGML_UNARY_OP_ELU:
case GGML_UNARY_OP_XIELU:
case GGML_UNARY_OP_CEIL:
return true;
case GGML_UNARY_OP_FLOOR:
+74 -45
View File
@@ -723,6 +723,7 @@ struct vk_device_struct {
bool uma;
bool prefer_host_memory;
bool float_controls_rte_fp16;
bool float_controls_denorm_preserve_fp16;
bool subgroup_basic;
bool subgroup_arithmetic;
bool subgroup_shuffle;
@@ -868,8 +869,9 @@ struct vk_device_struct {
vk_pipeline pipeline_cpy_f32_quant[GGML_TYPE_COUNT];
vk_pipeline pipeline_cpy_quant_f32[GGML_TYPE_COUNT];
vk_pipeline pipeline_cpy_transpose_16, pipeline_cpy_transpose_32;
vk_pipeline pipeline_set_rows_i32[GGML_TYPE_COUNT];
vk_pipeline pipeline_set_rows_i64[GGML_TYPE_COUNT];
// [src0 0=fp32,1=fp16][dst]
vk_pipeline pipeline_set_rows_i32[2][GGML_TYPE_COUNT];
vk_pipeline pipeline_set_rows_i64[2][GGML_TYPE_COUNT];
vk_pipeline pipeline_norm_f32;
vk_pipeline pipeline_group_norm_f32;
vk_pipeline pipeline_rms_norm_f32;
@@ -2595,10 +2597,10 @@ static void ggml_vk_create_pipeline_func(vk_device& device, vk_pipeline& pipelin
vk::ShaderModuleCreateInfo shader_module_create_info({}, spv_size, reinterpret_cast<const uint32_t *>(spv_data));
// Patch SPIR-V to enable RTE rounding for FP16, avoiding the need for
// separate shader variants compiled with -DRTE16.
// Patch SPIR-V to enable supported FP16 float controls, avoiding the need
// for separate shader variants.
std::vector<uint32_t> spirv;
if (device->float_controls_rte_fp16) {
if (device->float_controls_rte_fp16 || device->float_controls_denorm_preserve_fp16) {
const uint32_t* spv_words = reinterpret_cast<const uint32_t *>(spv_data);
size_t word_count = spv_size / sizeof(uint32_t);
spirv.assign(spv_words, spv_words + word_count);
@@ -2635,9 +2637,17 @@ static void ggml_vk_create_pipeline_func(vk_device& device, vk_pipeline& pipelin
// Insert from latest position first so earlier indices stay valid.
// OpExecutionMode %entrypoint RoundingModeRTE 16
uint32_t exec_mode[] = { (4u << spv::WordCountShift) | spv::OpExecutionMode, entry_point_id, spv::ExecutionModeRoundingModeRTE, 16 };
spirv.insert(spirv.begin() + exec_insert_pos, std::begin(exec_mode), std::end(exec_mode));
if (device->float_controls_rte_fp16) {
// OpExecutionMode %entrypoint RoundingModeRTE 16
uint32_t exec_mode[] = { (4u << spv::WordCountShift) | spv::OpExecutionMode, entry_point_id, spv::ExecutionModeRoundingModeRTE, 16 };
spirv.insert(spirv.begin() + exec_insert_pos, std::begin(exec_mode), std::end(exec_mode));
}
if (device->float_controls_denorm_preserve_fp16) {
// OpExecutionMode %entrypoint DenormPreserve 16
uint32_t exec_mode[] = { (4u << spv::WordCountShift) | spv::OpExecutionMode, entry_point_id, spv::ExecutionModeDenormPreserve, 16 };
spirv.insert(spirv.begin() + exec_insert_pos, std::begin(exec_mode), std::end(exec_mode));
}
// OpExtension "SPV_KHR_float_controls"
const char ext_str[] = "SPV_KHR_float_controls";
@@ -2647,9 +2657,17 @@ static void ggml_vk_create_pipeline_func(vk_device& device, vk_pipeline& pipelin
memcpy(&extension[1], ext_str, sizeof(ext_str));
spirv.insert(spirv.begin() + ext_insert_pos, extension.begin(), extension.end());
// OpCapability RoundingModeRTE
uint32_t capability[] = { (2u << spv::WordCountShift) | spv::OpCapability, spv::CapabilityRoundingModeRTE };
spirv.insert(spirv.begin() + cap_insert_pos, std::begin(capability), std::end(capability));
if (device->float_controls_rte_fp16) {
// OpCapability RoundingModeRTE
uint32_t capability[] = { (2u << spv::WordCountShift) | spv::OpCapability, spv::CapabilityRoundingModeRTE };
spirv.insert(spirv.begin() + cap_insert_pos, std::begin(capability), std::end(capability));
}
if (device->float_controls_denorm_preserve_fp16) {
// OpCapability DenormPreserve
uint32_t capability[] = { (2u << spv::WordCountShift) | spv::OpCapability, spv::CapabilityDenormPreserve };
spirv.insert(spirv.begin() + cap_insert_pos, std::begin(capability), std::end(capability));
}
shader_module_create_info = vk::ShaderModuleCreateInfo({}, spirv.size() * sizeof(uint32_t), spirv.data());
}
@@ -5187,20 +5205,22 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
ggml_vk_create_pipeline(device, device->pipeline_cpy_f32_quant[GGML_TYPE_Q8_0], "cpy_f32_q8_0", cpy_f32_q8_0_len, cpy_f32_q8_0_data, "main", 2, sizeof(vk_op_unary_push_constants), {32, 1, 1}, {}, 1);
ggml_vk_create_pipeline(device, device->pipeline_cpy_f32_quant[GGML_TYPE_IQ4_NL], "cpy_f32_iq4_nl", cpy_f32_iq4_nl_len, cpy_f32_iq4_nl_data, "main", 2, sizeof(vk_op_unary_push_constants), {32, 1, 1}, {}, 1);
#define SET_ROWS(itype) \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [GGML_TYPE_F32], "set_rows_f32" #itype, set_rows_f32 ## itype ## _len, set_rows_f32 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [GGML_TYPE_F16], "set_rows_f16" #itype, set_rows_f16 ## itype ## _len, set_rows_f16 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [GGML_TYPE_BF16], "set_rows_bf16" #itype, set_rows_bf16 ## itype ## _len, set_rows_bf16 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [GGML_TYPE_Q1_0], "set_rows_q1_0" #itype, set_rows_q1_0 ## itype ## _len, set_rows_q1_0 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [GGML_TYPE_Q4_0], "set_rows_q4_0" #itype, set_rows_q4_0 ## itype ## _len, set_rows_q4_0 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [GGML_TYPE_Q4_1], "set_rows_q4_1" #itype, set_rows_q4_1 ## itype ## _len, set_rows_q4_1 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [GGML_TYPE_Q5_0], "set_rows_q5_0" #itype, set_rows_q5_0 ## itype ## _len, set_rows_q5_0 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [GGML_TYPE_Q5_1], "set_rows_q5_1" #itype, set_rows_q5_1 ## itype ## _len, set_rows_q5_1 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [GGML_TYPE_Q8_0], "set_rows_q8_0" #itype, set_rows_q8_0 ## itype ## _len, set_rows_q8_0 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [GGML_TYPE_IQ4_NL], "set_rows_iq4_nl" #itype, set_rows_iq4_nl ## itype ## _len, set_rows_iq4_nl ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true);
#define SET_ROWS(src_idx, src, itype) \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [src_idx][GGML_TYPE_F32], "set_rows_" #src "_f32" #itype, set_rows_ ## src ## _f32 ## itype ## _len, set_rows_ ## src ## _f32 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [src_idx][GGML_TYPE_F16], "set_rows_" #src "_f16" #itype, set_rows_ ## src ## _f16 ## itype ## _len, set_rows_ ## src ## _f16 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [src_idx][GGML_TYPE_BF16], "set_rows_" #src "_bf16" #itype, set_rows_ ## src ## _bf16 ## itype ## _len, set_rows_ ## src ## _bf16 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [src_idx][GGML_TYPE_Q1_0], "set_rows_" #src "_q1_0" #itype, set_rows_ ## src ## _q1_0 ## itype ## _len, set_rows_ ## src ## _q1_0 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [src_idx][GGML_TYPE_Q4_0], "set_rows_" #src "_q4_0" #itype, set_rows_ ## src ## _q4_0 ## itype ## _len, set_rows_ ## src ## _q4_0 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [src_idx][GGML_TYPE_Q4_1], "set_rows_" #src "_q4_1" #itype, set_rows_ ## src ## _q4_1 ## itype ## _len, set_rows_ ## src ## _q4_1 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [src_idx][GGML_TYPE_Q5_0], "set_rows_" #src "_q5_0" #itype, set_rows_ ## src ## _q5_0 ## itype ## _len, set_rows_ ## src ## _q5_0 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [src_idx][GGML_TYPE_Q5_1], "set_rows_" #src "_q5_1" #itype, set_rows_ ## src ## _q5_1 ## itype ## _len, set_rows_ ## src ## _q5_1 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [src_idx][GGML_TYPE_Q8_0], "set_rows_" #src "_q8_0" #itype, set_rows_ ## src ## _q8_0 ## itype ## _len, set_rows_ ## src ## _q8_0 ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true); \
ggml_vk_create_pipeline(device, device->pipeline_set_rows ## itype [src_idx][GGML_TYPE_IQ4_NL], "set_rows_" #src "_iq4_nl" #itype, set_rows_ ## src ## _iq4_nl ## itype ## _len, set_rows_ ## src ## _iq4_nl ## itype ## _data, "main", 3, sizeof(vk_op_binary_push_constants), {1, 1, 1}, {1}, 1, true);
SET_ROWS(_i32)
SET_ROWS(_i64)
SET_ROWS(0, f32, _i32)
SET_ROWS(0, f32, _i64)
SET_ROWS(1, f16, _i32)
SET_ROWS(1, f16, _i64)
#undef SET_ROWS
@@ -6031,6 +6051,7 @@ static vk_device ggml_vk_get_device(size_t idx) {
device->shader_core_count = 0;
}
device->float_controls_rte_fp16 = vk12_props.shaderRoundingModeRTEFloat16;
device->float_controls_denorm_preserve_fp16 = vk12_props.shaderDenormPreserveFloat16;
device->subgroup_basic = (vk11_props.subgroupSupportedStages & vk::ShaderStageFlagBits::eCompute) &&
(vk11_props.subgroupSupportedOperations & vk::SubgroupFeatureFlagBits::eBasic);
@@ -10843,10 +10864,17 @@ static vk_pipeline ggml_vk_op_get_pipeline(ggml_backend_vk_context * ctx, const
case GGML_OP_DUP:
return ggml_vk_get_cpy_pipeline(ctx, src0, dst, dst->type);
case GGML_OP_SET_ROWS:
if (src1->type == GGML_TYPE_I64) {
return ctx->device->pipeline_set_rows_i64[dst->type];
} else {
return ctx->device->pipeline_set_rows_i32[dst->type];
{
if (src0->type != GGML_TYPE_F32 && src0->type != GGML_TYPE_F16) {
return nullptr;
}
const int src_idx = src0->type == GGML_TYPE_F16;
if (src1->type == GGML_TYPE_I64) {
return ctx->device->pipeline_set_rows_i64[src_idx][dst->type];
} else if (src1->type == GGML_TYPE_I32) {
return ctx->device->pipeline_set_rows_i32[src_idx][dst->type];
}
return nullptr;
}
case GGML_OP_SILU_BACK:
if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
@@ -17500,24 +17528,25 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
return op->type == GGML_TYPE_F32 && op->src[0]->type == GGML_TYPE_F32;
case GGML_OP_SET_ROWS:
{
if (op->src[0]->type == GGML_TYPE_F32) {
switch (op->type) {
case GGML_TYPE_F32:
case GGML_TYPE_F16:
case GGML_TYPE_BF16:
case GGML_TYPE_Q1_0:
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q4_1:
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1:
case GGML_TYPE_Q8_0:
case GGML_TYPE_IQ4_NL:
return true;
default:
return false;
}
if ((op->src[0]->type != GGML_TYPE_F32 && op->src[0]->type != GGML_TYPE_F16) ||
(op->src[1]->type != GGML_TYPE_I32 && op->src[1]->type != GGML_TYPE_I64)) {
return false;
}
switch (op->type) {
case GGML_TYPE_F32:
case GGML_TYPE_F16:
case GGML_TYPE_BF16:
case GGML_TYPE_Q1_0:
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q4_1:
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1:
case GGML_TYPE_Q8_0:
case GGML_TYPE_IQ4_NL:
return true;
default:
return false;
}
return false;
}
case GGML_OP_CONT:
case GGML_OP_CPY:
@@ -10,7 +10,7 @@ layout(local_size_x = 32, local_size_y = 1, local_size_z = 1) in;
const uint BLOCK_SIZE = 32;
#endif
layout (binding = 0) readonly buffer S {float data_s[];};
layout (binding = 0) readonly buffer S {S_TYPE data_s[];};
#if defined(SET_ROWS)
#include "generic_binary_head.glsl"
@@ -35,7 +35,7 @@ void quantize(uint dst_idx, uint src_idx)
float vmax = 0.0;
[[unroll]] for (int j = 0; j < QUANT_K_Q4_0; ++j) {
const float v = data_s[src_idx + j];
const float v = float(data_s[src_idx + j]);
if (amax < abs(v)) {
amax = abs(v);
vmax = v;
@@ -48,8 +48,8 @@ void quantize(uint dst_idx, uint src_idx)
data_q[dst_idx].d = float16_t(d);
[[unroll]] for (int j = 0; j < QUANT_K_Q4_0/2; ++j) {
const float x0 = data_s[src_idx + 0 + j]*id;
const float x1 = data_s[src_idx + QUANT_K_Q4_0/2 + j]*id;
const float x0 = float(data_s[src_idx + 0 + j])*id;
const float x1 = float(data_s[src_idx + QUANT_K_Q4_0/2 + j])*id;
const uint xi0 = min(15, int(x0 + 8.5));
const uint xi1 = min(15, int(x1 + 8.5));
@@ -66,7 +66,7 @@ void quantize(uint dst_idx, uint src_idx)
float vmax = -vmin;
[[unroll]] for (int j = 0; j < QUANT_K_Q4_1; ++j) {
const float v = data_s[src_idx + j];
const float v = float(data_s[src_idx + j]);
if (v < vmin) vmin = v;
if (v > vmax) vmax = v;
@@ -79,8 +79,8 @@ void quantize(uint dst_idx, uint src_idx)
data_q[dst_idx].m = float16_t(vmin);
[[unroll]] for (int j = 0; j < QUANT_K_Q4_1/2; ++j) {
const float x0 = (data_s[src_idx + 0 + j] - vmin)*id;
const float x1 = (data_s[src_idx + QUANT_K_Q4_1/2 + j] - vmin)*id;
const float x0 = (float(data_s[src_idx + 0 + j]) - vmin)*id;
const float x1 = (float(data_s[src_idx + QUANT_K_Q4_1/2 + j]) - vmin)*id;
const uint xi0 = min(15, int(x0 + 0.5));
const uint xi1 = min(15, int(x1 + 0.5));
@@ -97,7 +97,7 @@ void quantize(uint dst_idx, uint src_idx)
float vmax = 0.0;
[[unroll]] for (int j = 0; j < QUANT_K_Q5_0; ++j) {
const float v = data_s[src_idx + j];
const float v = float(data_s[src_idx + j]);
if (amax < abs(v)) {
amax = abs(v);
vmax = v;
@@ -111,8 +111,8 @@ void quantize(uint dst_idx, uint src_idx)
uint32_t qh = 0;
[[unroll]] for (int j = 0; j < QUANT_K_Q5_0/2; ++j) {
const float x0 = data_s[src_idx + 0 + j]*id;
const float x1 = data_s[src_idx + QUANT_K_Q5_0/2 + j]*id;
const float x0 = float(data_s[src_idx + 0 + j])*id;
const float x1 = float(data_s[src_idx + QUANT_K_Q5_0/2 + j])*id;
const uint xi0 = min(31, int(x0 + 16.5));
const uint xi1 = min(31, int(x1 + 16.5));
@@ -129,11 +129,11 @@ void quantize(uint dst_idx, uint src_idx)
#if defined(DATA_A_Q5_1)
void quantize(uint dst_idx, uint src_idx)
{
float min = data_s[src_idx + 0];
float min = float(data_s[src_idx + 0]);
float max = min;
[[unroll]] for (int j = 1; j < QUANT_K_Q5_1; ++j) {
const float v = data_s[src_idx + j];
const float v = float(data_s[src_idx + j]);
min = v < min ? v : min;
max = v > max ? v : max;
}
@@ -146,8 +146,8 @@ void quantize(uint dst_idx, uint src_idx)
uint32_t qh = 0;
[[unroll]] for (int j = 0; j < QUANT_K_Q5_1/2; ++j) {
const float x0 = (data_s[src_idx + 0 + j] - min)*id;
const float x1 = (data_s[src_idx + QUANT_K_Q5_1/2 + j] - min)*id;
const float x0 = (float(data_s[src_idx + 0 + j]) - min)*id;
const float x1 = (float(data_s[src_idx + QUANT_K_Q5_1/2 + j]) - min)*id;
const uint xi0 = uint(x0 + 0.5);
const uint xi1 = uint(x1 + 0.5);
@@ -166,7 +166,7 @@ void quantize(uint dst_idx, uint src_idx)
float amax = 0.0; // absolute max
[[unroll]] for (int j = 0; j < QUANT_K_Q8_0; j++) {
const float v = data_s[src_idx + j];
const float v = float(data_s[src_idx + j]);
amax = max(amax, abs(v));
}
@@ -176,7 +176,7 @@ void quantize(uint dst_idx, uint src_idx)
data_q[dst_idx].d = float16_t(d);
[[unroll]] for (int j = 0; j < QUANT_K_Q8_0; ++j) {
const float x0 = data_s[src_idx + j]*id;
const float x0 = float(data_s[src_idx + j])*id;
data_q[dst_idx].qs[j] = int8_t(round(x0));
}
@@ -189,7 +189,7 @@ void quantize(uint dst_idx, uint src_idx)
float sum_abs = 0.0;
[[unroll]] for (int j = 0; j < QUANT_K_Q1_0; j++) {
sum_abs += abs(data_s[src_idx + j]);
sum_abs += abs(float(data_s[src_idx + j]));
}
const float d = sum_abs / QUANT_K_Q1_0;
@@ -201,7 +201,7 @@ void quantize(uint dst_idx, uint src_idx)
}
[[unroll]] for (int j = 0; j < QUANT_K_Q1_0; ++j) {
if (data_s[src_idx + j] >= 0.0) {
if (float(data_s[src_idx + j]) >= 0.0) {
data_q[dst_idx].qs[j / 8] |= uint8_t(1 << (j % 8));
}
}
@@ -226,7 +226,7 @@ void quantize(uint dst_idx, uint src_idx)
float vmax = 0.0;
[[unroll]] for (int j = 0; j < QUANT_K_IQ4_NL; ++j) {
const float v = data_s[src_idx + j];
const float v = float(data_s[src_idx + j]);
if (amax < abs(v)) {
amax = abs(v);
vmax = v;
@@ -238,16 +238,16 @@ void quantize(uint dst_idx, uint src_idx)
float sumqx = 0, sumq2 = 0;
[[unroll]] for (int j = 0; j < QUANT_K_IQ4_NL/2; ++j) {
const float x0 = data_s[src_idx + 0 + j]*id;
const float x1 = data_s[src_idx + QUANT_K_IQ4_NL/2 + j]*id;
const float x0 = float(data_s[src_idx + 0 + j])*id;
const float x1 = float(data_s[src_idx + QUANT_K_IQ4_NL/2 + j])*id;
const uint xi0 = best_index(x0);
const uint xi1 = best_index(x1);
data_q[dst_idx].qs[j] = uint8_t(xi0 | (xi1 << 4));
const float v0 = kvalues_iq4nl[xi0];
const float v1 = kvalues_iq4nl[xi1];
const float w0 = data_s[src_idx + 0 + j]*data_s[src_idx + 0 + j];
const float w1 = data_s[src_idx + QUANT_K_IQ4_NL/2 + j]*data_s[src_idx + QUANT_K_IQ4_NL/2 + j];
sumqx += w0*v0*data_s[src_idx + j] + w1*v1*data_s[src_idx + QUANT_K_IQ4_NL/2 + j];
const float w0 = float(data_s[src_idx + 0 + j])*float(data_s[src_idx + 0 + j]);
const float w1 = float(data_s[src_idx + QUANT_K_IQ4_NL/2 + j])*float(data_s[src_idx + QUANT_K_IQ4_NL/2 + j]);
sumqx += w0*v0*float(data_s[src_idx + j]) + w1*v1*float(data_s[src_idx + QUANT_K_IQ4_NL/2 + j]);
sumq2 += w0*v0*v0 + w1*v1*v1;
}
@@ -259,14 +259,14 @@ void quantize(uint dst_idx, uint src_idx)
#if defined(DATA_A_F32) || defined(DATA_A_F16)
void quantize(uint dst_idx, uint src_idx)
{
data_q[dst_idx] = A_TYPE(data_s[src_idx]);
data_q[dst_idx] = A_TYPE(float(data_s[src_idx]));
}
#endif
#if defined(DATA_A_BF16)
void quantize(uint dst_idx, uint src_idx)
{
data_q[dst_idx] = A_TYPE(fp32_to_bf16(data_s[src_idx]));
data_q[dst_idx] = A_TYPE(fp32_to_bf16(float(data_s[src_idx])));
}
#endif
@@ -824,13 +824,15 @@ void process_shaders() {
string_to_spv("cpy_transpose_32", "copy_transpose.comp", {{"A_TYPE", "uint"}, {"D_TYPE", "uint"}});
for (std::string t : {"q1_0", "q4_0", "q4_1", "q5_0", "q5_1", "q8_0", "iq4_nl"}) {
string_to_spv("cpy_f32_" + t, "copy_to_quant.comp", {{"DATA_A_" + to_uppercase(t), "1"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}});
string_to_spv("cpy_f32_" + t, "copy_to_quant.comp", {{"DATA_A_" + to_uppercase(t), "1"}, {"S_TYPE", "float"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}});
string_to_spv("cpy_" + t + "_f32", "copy_from_quant.comp", {{"DATA_A_" + to_uppercase(t), "1"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}});
}
for (std::string t : {"f32", "f16", "bf16", "q1_0", "q4_0", "q4_1", "q5_0", "q5_1", "q8_0", "iq4_nl"}) {
string_to_spv("set_rows_" + t + "_i32", "copy_to_quant.comp", {{"SET_ROWS", "1"}, {"DATA_A_" + to_uppercase(t), "1"}, {"B_TYPE", "uint"}, {"B_SIZE", "32"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}});
string_to_spv("set_rows_" + t + "_i64", "copy_to_quant.comp", {{"SET_ROWS", "1"}, {"DATA_A_" + to_uppercase(t), "1"}, {"B_TYPE", "uvec2"}, {"B_SIZE", "64"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}});
for (auto src : {std::pair{"f32", "float"}, std::pair{"f16", "float16_t"}}) {
for (std::string dst : {"f32", "f16", "bf16", "q1_0", "q4_0", "q4_1", "q5_0", "q5_1", "q8_0", "iq4_nl"}) {
string_to_spv("set_rows_" + std::string(src.first) + "_" + dst + "_i32", "copy_to_quant.comp", {{"SET_ROWS", "1"}, {"DATA_A_" + to_uppercase(dst), "1"}, {"B_TYPE", "uint"}, {"B_SIZE", "32"}, {"S_TYPE", src.second}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}});
string_to_spv("set_rows_" + std::string(src.first) + "_" + dst + "_i64", "copy_to_quant.comp", {{"SET_ROWS", "1"}, {"DATA_A_" + to_uppercase(dst), "1"}, {"B_TYPE", "uvec2"}, {"B_SIZE", "64"}, {"S_TYPE", src.second}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}});
}
}
auto get_type_str = [](bool f16) {
+23 -11
View File
@@ -1464,14 +1464,14 @@ bool ggml_is_transposed(const struct ggml_tensor * tensor) {
return tensor->nb[0] > tensor->nb[1];
}
static bool ggml_is_contiguous_n(const struct ggml_tensor * tensor, int n) {
static bool ggml_is_contiguous_m_n(const struct ggml_tensor * tensor, int m, int n) {
size_t next_nb = ggml_type_size(tensor->type);
if (tensor->ne[0] != ggml_blck_size(tensor->type) && tensor->nb[0] != next_nb) {
return false;
}
next_nb *= tensor->ne[0]/ggml_blck_size(tensor->type);
for (int i = 1; i < GGML_MAX_DIMS; i++) {
if (i > n) {
for (int i = 1; i < n; i++) {
if (i > m) {
if (tensor->ne[i] != 1 && tensor->nb[i] != next_nb) {
return false;
}
@@ -1489,15 +1489,27 @@ bool ggml_is_contiguous(const struct ggml_tensor * tensor) {
}
bool ggml_is_contiguous_0(const struct ggml_tensor * tensor) {
return ggml_is_contiguous_n(tensor, 0);
return ggml_is_contiguous_m_n(tensor, 0, GGML_MAX_DIMS);
}
bool ggml_is_contiguous_1(const struct ggml_tensor * tensor) {
return ggml_is_contiguous_n(tensor, 1);
return ggml_is_contiguous_m_n(tensor, 1, GGML_MAX_DIMS);
}
bool ggml_is_contiguous_2(const struct ggml_tensor * tensor) {
return ggml_is_contiguous_n(tensor, 2);
return ggml_is_contiguous_m_n(tensor, 2, GGML_MAX_DIMS);
}
bool ggml_is_contiguous_to_1(const struct ggml_tensor * tensor) {
return ggml_is_contiguous_m_n(tensor, 0, 1);
}
bool ggml_is_contiguous_to_2(const struct ggml_tensor * tensor) {
return ggml_is_contiguous_m_n(tensor, 0, 2);
}
bool ggml_is_contiguous_to_3(const struct ggml_tensor * tensor) {
return ggml_is_contiguous_m_n(tensor, 0, 3);
}
bool ggml_is_contiguously_allocated(const struct ggml_tensor * tensor) {
@@ -4507,7 +4519,7 @@ struct ggml_tensor * ggml_conv_1d(
int s0,
int p0,
int d0) {
struct ggml_tensor * im2col = ggml_im2col(ctx, a, b, s0, 0, p0, 0, d0, 0, false, GGML_TYPE_F16); // [N, OL, IC * K]
struct ggml_tensor * im2col = ggml_im2col(ctx, a, b, s0, 0, p0, 0, d0, 0, false, a->type == GGML_TYPE_BF16 ? GGML_TYPE_F32 : GGML_TYPE_F16); // [N, OL, IC * K]
struct ggml_tensor * result =
ggml_mul_mat(ctx,
@@ -4541,7 +4553,7 @@ struct ggml_tensor * ggml_conv_1d_dw(
int d0) {
struct ggml_tensor * new_b = ggml_reshape_4d(ctx, b, b->ne[0], 1, b->ne[1], b->ne[2]);
struct ggml_tensor * im2col = ggml_im2col(ctx, a, new_b, s0, 0, p0, 0, d0, 0, false, GGML_TYPE_F16);
struct ggml_tensor * im2col = ggml_im2col(ctx, a, new_b, s0, 0, p0, 0, d0, 0, false, a->type == GGML_TYPE_BF16 ? GGML_TYPE_F32 : GGML_TYPE_F16);
struct ggml_tensor * result = ggml_mul_mat(ctx, im2col, a);
@@ -4647,7 +4659,7 @@ struct ggml_tensor * ggml_conv_2d(
int p1,
int d0,
int d1) {
struct ggml_tensor * im2col = ggml_im2col(ctx, a, b, s0, s1, p0, p1, d0, d1, true, a->type); // [N, OH, OW, IC * KH * KW]
struct ggml_tensor * im2col = ggml_im2col(ctx, a, b, s0, s1, p0, p1, d0, d1, true, a->type == GGML_TYPE_BF16 ? GGML_TYPE_F32 : GGML_TYPE_F16); // [N, OH, OW, IC * KH * KW]
struct ggml_tensor * result =
ggml_mul_mat(ctx,
@@ -4729,7 +4741,7 @@ struct ggml_tensor * ggml_conv_3d(
int d1, // dilation height
int d2 // dilation depth
) {
struct ggml_tensor * im2col = ggml_im2col_3d(ctx, a, b, IC, s0, s1, s2, p0, p1, p2, d0, d1, d2, a->type); // [N*OD, OH, OW, IC * KD * KH * KW]
struct ggml_tensor * im2col = ggml_im2col_3d(ctx, a, b, IC, s0, s1, s2, p0, p1, p2, d0, d1, d2, a->type == GGML_TYPE_BF16 ? GGML_TYPE_F32 : GGML_TYPE_F16); // [N*OD, OH, OW, IC * KD * KH * KW]
int64_t OC = a->ne[3] / IC;
int64_t N = b->ne[3] / IC;
@@ -4779,7 +4791,7 @@ struct ggml_tensor * ggml_conv_2d_dw(
struct ggml_tensor * new_a = ggml_reshape_4d(ctx, a, a->ne[0], a->ne[1], 1, a->ne[2] * a->ne[3]);
struct ggml_tensor * im2col = ggml_im2col(ctx, new_a,
ggml_reshape_4d(ctx, b, b->ne[0], b->ne[1], 1, b->ne[2] * b->ne[3]),
s0, s1, p0, p1, d0, d1, true, GGML_TYPE_F16); // [N * IC, OH, OW, KH * KW]
s0, s1, p0, p1, d0, d1, true, a->type == GGML_TYPE_BF16 ? GGML_TYPE_F32 : GGML_TYPE_F16); // [N * IC, OH, OW, KH * KW]
struct ggml_tensor * new_b = ggml_reshape_4d(ctx, im2col, im2col->ne[0], im2col->ne[2] * im2col->ne[1], b->ne[2], b->ne[3]); // [N * IC, OH, OW, KH * KW] => [N, IC, OH * OW, KH * KW]
new_a = ggml_reshape_4d(ctx, new_a, (new_a->ne[0] * new_a->ne[1]), new_a->ne[2], new_a->ne[3], 1); // [OC1, KH, KW] => [1, OC, 1, KH * KW]
+85 -15
View File
@@ -720,7 +720,7 @@ llama_dsv4_comp_state::llama_dsv4_comp_state(
auto it = ctx_map.find(buft);
if (it == ctx_map.end()) {
ggml_init_params params = {
/*.mem_size =*/ size_t(2u*hparams.n_layer()*ggml_tensor_overhead()),
/*.mem_size =*/ size_t(2u*(1 + n_stream)*hparams.n_layer()*ggml_tensor_overhead()),
/*.mem_buffer =*/ NULL,
/*.no_alloc =*/ true,
};
@@ -767,9 +767,17 @@ llama_dsv4_comp_state::llama_dsv4_comp_state(
ggml_format_name(kv, "dsv4_%s_state_kv_l%d", name, il);
ggml_format_name(score, "dsv4_%s_state_score_l%d", name, il);
std::vector<ggml_tensor *> kv_stream;
std::vector<ggml_tensor *> score_stream;
for (uint32_t s = 0; s < n_stream; ++s) {
kv_stream.push_back(ggml_view_2d(ctx, kv, n_embd_state, state_size, kv->nb[1], s*kv->nb[2]));
score_stream.push_back(ggml_view_2d(ctx, score, n_embd_state, state_size, score->nb[1], s*score->nb[2]));
}
map_layer_ids[il] = layers.size();
layers.push_back({ il, kv, score });
layers.push_back({ il, kv, score, std::move(kv_stream), std::move(score_stream) });
}
for (auto & [buft, ctx] : ctx_map) {
@@ -809,6 +817,30 @@ void llama_dsv4_comp_state::clear(llama_seq_id seq_id, bool data) {
}
}
void llama_dsv4_comp_state::seq_cp(llama_seq_id seq_id_src, llama_seq_id seq_id_dst) {
GGML_ASSERT(seq_id_src >= 0 && (uint32_t) seq_id_src < n_stream);
GGML_ASSERT(seq_id_dst >= 0 && (uint32_t) seq_id_dst < n_stream);
if (seq_id_src == seq_id_dst) {
return;
}
sc_info.ssrc.push_back((uint32_t) seq_id_src);
sc_info.sdst.push_back((uint32_t) seq_id_dst);
}
void llama_dsv4_comp_state::apply_copies(const stream_copy_info & sc_info) const {
for (size_t i = 0; i < sc_info.ssrc.size(); ++i) {
const uint32_t ssrc = sc_info.ssrc[i];
const uint32_t sdst = sc_info.sdst[i];
for (const auto & layer : layers) {
ggml_backend_tensor_copy(layer.kv_stream[ssrc], layer.kv_stream[sdst]);
ggml_backend_tensor_copy(layer.score_stream[ssrc], layer.score_stream[sdst]);
}
}
}
uint32_t llama_dsv4_comp_state::get_ratio() const {
return ratio;
}
@@ -1154,7 +1186,13 @@ llama_memory_context_ptr llama_kv_cache_dsv4::init_full() {
}
llama_memory_context_ptr llama_kv_cache_dsv4::init_update(llama_context * lctx, bool optimize) {
return std::make_unique<llama_kv_cache_dsv4_context>(this, lctx, optimize);
return std::make_unique<llama_kv_cache_dsv4_context>(
this,
lctx,
optimize,
std::move(csa_state->sc_info),
std::move(hca_state->sc_info),
std::move(lid_state->sc_info));
}
bool llama_kv_cache_dsv4::get_can_shift() const {
@@ -1174,14 +1212,19 @@ bool llama_kv_cache_dsv4::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_pos p1
}
if (p0 > 0) {
// DSV4 compressed cache rows are derived from running compressor state,
// so arbitrary rollback is not reconstructible from the raw cache alone.
// Allow the common prompt-cache cleanup no-op: remove [end, infinity).
if (seq_id >= 0 && p0 > kv_raw->seq_pos_max(seq_id)) {
return true;
if (seq_id < 0 || (uint32_t) seq_id >= n_seq_max ||
p0 <= kv_raw->seq_pos_max(seq_id)) {
return false;
}
return false;
bool res = true;
res = res & kv_raw->seq_rm(seq_id, p0, -1);
res = res & kv_csa->seq_rm(seq_id, p0/DSV4_CSA_RATIO, -1);
res = res & kv_hca->seq_rm(seq_id, p0/DSV4_HCA_RATIO, -1);
res = res & kv_lid->seq_rm(seq_id, p0/DSV4_CSA_RATIO, -1);
return res;
}
const bool res = kv_raw->seq_rm(seq_id, p0, p1);
@@ -1194,7 +1237,16 @@ bool llama_kv_cache_dsv4::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_pos p1
}
void llama_kv_cache_dsv4::seq_cp(llama_seq_id seq_id_src, llama_seq_id seq_id_dst, llama_pos p0, llama_pos p1) {
GGML_ASSERT(p0 <= 0 && p1 < 0 && "DSV4 only supports full sequence copies");
kv_raw->seq_cp(seq_id_src, seq_id_dst, p0, p1);
kv_csa->seq_cp(seq_id_src, seq_id_dst, -1, -1);
kv_hca->seq_cp(seq_id_src, seq_id_dst, -1, -1);
kv_lid->seq_cp(seq_id_src, seq_id_dst, -1, -1);
csa_state->seq_cp(seq_id_src, seq_id_dst);
hca_state->seq_cp(seq_id_src, seq_id_dst);
lid_state->seq_cp(seq_id_src, seq_id_dst);
}
void llama_kv_cache_dsv4::seq_keep(llama_seq_id seq_id) {
@@ -1639,20 +1691,26 @@ llama_kv_cache_dsv4_context::llama_kv_cache_dsv4_context(
llama_kv_cache_dsv4_context::llama_kv_cache_dsv4_context(
llama_kv_cache_dsv4 * kv,
llama_context * lctx,
bool optimize) :
bool optimize,
stream_copy_info sc_info_csa,
stream_copy_info sc_info_hca,
stream_copy_info sc_info_lid) :
ctx_raw(std::make_unique<llama_kv_cache_dsv4_raw_context>(kv->get_raw(), lctx, optimize)),
ctx_csa_mem(kv->get_csa()->init_update(lctx, optimize)),
ctx_hca_mem(kv->get_hca()->init_update(lctx, optimize)),
ctx_lid_mem(kv->get_lid()->init_update(lctx, optimize)),
ctx_csa(std::make_unique<llama_kv_cache_dsv4_comp_context>(kv->get_csa())),
ctx_hca(std::make_unique<llama_kv_cache_dsv4_comp_context>(kv->get_hca())),
ctx_lid(std::make_unique<llama_kv_cache_dsv4_comp_context>(kv->get_lid())),
csa_state(kv->get_csa_state()),
hca_state(kv->get_hca_state()),
lid_state(kv->get_lid_state()),
sc_info_csa(std::move(sc_info_csa)),
sc_info_hca(std::move(sc_info_hca)),
sc_info_lid(std::move(sc_info_lid)),
status(llama_memory_status_combine(
llama_memory_status_combine(ctx_raw->get_status(), ctx_csa_mem->get_status()),
llama_memory_status_combine(ctx_hca_mem->get_status(), ctx_lid_mem->get_status()))) {
llama_memory_status_combine(
llama_memory_status_combine(ctx_raw->get_status(), ctx_csa_mem->get_status()),
llama_memory_status_combine(ctx_hca_mem->get_status(), ctx_lid_mem->get_status())),
this->sc_info_csa.empty() && this->sc_info_hca.empty() && this->sc_info_lid.empty() ?
LLAMA_MEMORY_STATUS_NO_UPDATE : LLAMA_MEMORY_STATUS_SUCCESS)) {
}
llama_kv_cache_dsv4_context::llama_kv_cache_dsv4_context(
@@ -1720,6 +1778,18 @@ bool llama_kv_cache_dsv4_context::apply() {
res = res & ctx_raw->apply();
if (ctx_csa_mem) {
res = res & ctx_csa_mem->apply();
res = res & ctx_hca_mem->apply();
res = res & ctx_lid_mem->apply();
}
if (ubatches.empty()) {
csa_state->apply_copies(sc_info_csa);
hca_state->apply_copies(sc_info_hca);
lid_state->apply_copies(sc_info_lid);
}
return res;
}
+21 -4
View File
@@ -10,6 +10,10 @@
class llama_dsv4_comp_state {
public:
using stream_copy_info = llama_kv_cache::stream_copy_info;
stream_copy_info sc_info;
llama_dsv4_comp_state(
const llama_model & model,
bool offload,
@@ -22,6 +26,8 @@ public:
const llama_memory_i::layer_filter_cb & filter);
void clear(llama_seq_id seq_id, bool data);
void seq_cp(llama_seq_id seq_id_src, llama_seq_id seq_id_dst);
void apply_copies(const stream_copy_info & sc_info) const;
uint32_t get_ratio() const;
uint32_t get_state_size() const;
@@ -44,6 +50,9 @@ private:
ggml_tensor * kv;
ggml_tensor * score;
std::vector<ggml_tensor *> kv_stream;
std::vector<ggml_tensor *> score_stream;
};
const uint32_t ratio;
@@ -245,6 +254,7 @@ private:
class llama_kv_cache_dsv4_context : public llama_memory_context_i {
public:
using slot_info_vec_t = llama_kv_cache::slot_info_vec_t;
using stream_copy_info = llama_kv_cache::stream_copy_info;
struct comp_plan {
// Per-ubatch recipe for updating compressor state, committing completed
@@ -291,7 +301,10 @@ public:
llama_kv_cache_dsv4_context(
llama_kv_cache_dsv4 * kv,
llama_context * lctx,
bool optimize);
bool optimize,
stream_copy_info sc_info_csa,
stream_copy_info sc_info_hca,
stream_copy_info sc_info_lid);
llama_kv_cache_dsv4_context(
llama_kv_cache_dsv4 * kv,
@@ -351,9 +364,13 @@ private:
const std::unique_ptr<llama_kv_cache_dsv4_comp_context> ctx_hca;
const std::unique_ptr<llama_kv_cache_dsv4_comp_context> ctx_lid;
const llama_dsv4_comp_state * csa_state = nullptr;
const llama_dsv4_comp_state * hca_state = nullptr;
const llama_dsv4_comp_state * lid_state = nullptr;
llama_dsv4_comp_state * csa_state = nullptr;
llama_dsv4_comp_state * hca_state = nullptr;
llama_dsv4_comp_state * lid_state = nullptr;
stream_copy_info sc_info_csa;
stream_copy_info sc_info_hca;
stream_copy_info sc_info_lid;
bool reserve_plans = false;
mutable comp_plan reserve_plan_csa;
+14 -12
View File
@@ -435,27 +435,29 @@ ggml_tensor * llama_model_deepseek4::graph::build_overlap_compressed_kv_from_sta
kv_state = dsv4_append_zero_row(ctx0, kv_state, false);
score_state = dsv4_append_zero_row(ctx0, score_state, true);
ggml_tensor * prev_idxs = dsv4_view_1d(ctx0, state_read_idxs, ratio*n_blocks, 0);
ggml_tensor * cur_idxs = dsv4_view_1d(ctx0, state_read_idxs, ratio*n_blocks, ratio*n_blocks);
const int64_t n_read = ratio*n_blocks;
ggml_tensor * kv_prev = ggml_get_rows(ctx0, kv_state, prev_idxs);
kv_prev = ggml_cont(ctx0, ggml_view_2d(ctx0, kv_prev, n_embd_head, ratio*n_blocks, kv_prev->nb[1], 0));
ggml_tensor * kv_rows = ggml_get_rows(ctx0, kv_state, state_read_idxs);
ggml_tensor * score_rows = ggml_get_rows(ctx0, score_state, state_read_idxs);
ggml_tensor * kv_prev = ggml_cont(ctx0,
ggml_view_2d(ctx0, kv_rows, n_embd_head, n_read, kv_rows->nb[1], 0));
kv_prev = ggml_reshape_3d(ctx0, kv_prev, n_embd_head, ratio, n_blocks);
cb(kv_prev, name, il);
ggml_tensor * score_prev = ggml_get_rows(ctx0, score_state, prev_idxs);
score_prev = ggml_cont(ctx0, ggml_view_2d(ctx0, score_prev, n_embd_head, ratio*n_blocks, score_prev->nb[1], 0));
ggml_tensor * score_prev = ggml_cont(ctx0,
ggml_view_2d(ctx0, score_rows, n_embd_head, n_read, score_rows->nb[1], 0));
score_prev = ggml_reshape_3d(ctx0, score_prev, n_embd_head, ratio, n_blocks);
cb(score_prev, name, il);
ggml_tensor * kv_cur = ggml_get_rows(ctx0, kv_state, cur_idxs);
kv_cur = ggml_cont(ctx0, ggml_view_2d(ctx0, kv_cur, n_embd_head, ratio*n_blocks, kv_cur->nb[1],
ggml_row_size(kv_cur->type, n_embd_head)));
ggml_tensor * kv_cur = ggml_cont(ctx0,
ggml_view_2d(ctx0, kv_rows, n_embd_head, n_read, kv_rows->nb[1],
n_read*kv_rows->nb[1] + ggml_row_size(kv_rows->type, n_embd_head)));
kv_cur = ggml_reshape_3d(ctx0, kv_cur, n_embd_head, ratio, n_blocks);
ggml_tensor * score_cur = ggml_get_rows(ctx0, score_state, cur_idxs);
score_cur = ggml_cont(ctx0, ggml_view_2d(ctx0, score_cur, n_embd_head, ratio*n_blocks, score_cur->nb[1],
ggml_row_size(score_cur->type, n_embd_head)));
ggml_tensor * score_cur = ggml_cont(ctx0,
ggml_view_2d(ctx0, score_rows, n_embd_head, n_read, score_rows->nb[1],
n_read*score_rows->nb[1] + ggml_row_size(score_rows->type, n_embd_head)));
score_cur = ggml_reshape_3d(ctx0, score_cur, n_embd_head, ratio, n_blocks);
ggml_tensor * values = ggml_concat(ctx0, kv_prev, kv_cur, 1);
+18 -13
View File
@@ -2423,11 +2423,10 @@ struct test_set_rows : public test_case {
void initialize_tensors(ggml_context * ctx) override {
for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
if (ggml_is_view_op(t->op)) {
continue;
}
if (t->type == GGML_TYPE_I64 || t->type == GGML_TYPE_I32) {
if (ggml_is_view_op(t->op)) {
continue;
}
init_set_rows_row_ids(t, ne[1]);
} else {
init_tensor_uniform(t);
@@ -2450,6 +2449,9 @@ struct test_set_rows : public test_case {
err_estimate /= 8.0f;
}
err_estimate *= err_estimate;
if (type_src == GGML_TYPE_F16) {
err_estimate *= 16.0f;
}
err_estimate /= 0.25f*float(ne[0] * r * ne[2]*nr23[0] * ne[3]*nr23[1]);
return err_estimate;
}
@@ -7928,17 +7930,19 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
test_cases.emplace_back(new test_set_rows(GGML_TYPE_F32, GGML_TYPE_F32, GGML_TYPE_I64, { 1, 8, 1, 3 }, { 1, 1 }, 2, false));
test_cases.emplace_back(new test_set_rows(GGML_TYPE_F32, GGML_TYPE_F32, GGML_TYPE_I32, { 1, 8, 1, 3 }, { 1, 1 }, 2, false));
test_cases.emplace_back(new test_set_rows(GGML_TYPE_F32, GGML_TYPE_Q8_0, GGML_TYPE_I32, { 256, 5, 1, 3 }, { 1, 1, }, 1, false));
for (ggml_type type : all_types) {
for (int b : {1, 7}) {
for (bool v : {false, true}) {
test_cases.emplace_back(new test_set_rows(GGML_TYPE_F32, type, GGML_TYPE_I64, { 256, 5, b, 3 }, { 1, 1, }, 1, v));
test_cases.emplace_back(new test_set_rows(GGML_TYPE_F32, type, GGML_TYPE_I64, { 256, 11, 1, b }, { 2, 3, }, 7, v));
for (ggml_type src_type : {GGML_TYPE_F16, GGML_TYPE_F32}) {
for (ggml_type type : all_types) {
for (int b : {1, 7}) {
for (bool v : {false, true}) {
test_cases.emplace_back(new test_set_rows(src_type, type, GGML_TYPE_I64, { 256, 5, b, 3 }, { 1, 1, }, 1, v));
test_cases.emplace_back(new test_set_rows(src_type, type, GGML_TYPE_I64, { 256, 11, 1, b }, { 2, 3, }, 7, v));
test_cases.emplace_back(new test_set_rows(GGML_TYPE_F32, type, GGML_TYPE_I64, { 3*ggml_blck_size(type), 3, b, 1 }, { 2, 3, }, 2, v));
test_cases.emplace_back(new test_set_rows(src_type, type, GGML_TYPE_I64, { 3*ggml_blck_size(type), 3, b, 1 }, { 2, 3, }, 2, v));
if (ggml_blck_size(type) == 1) {
test_cases.emplace_back(new test_set_rows(GGML_TYPE_F32, type, GGML_TYPE_I64, { 31, 3, b, 1 }, { 2, 3, }, 2, v));
test_cases.emplace_back(new test_set_rows(GGML_TYPE_F32, type, GGML_TYPE_I64, { 33, 5, 1, b }, { 2, 3, }, 1, v));
if (ggml_blck_size(type) == 1) {
test_cases.emplace_back(new test_set_rows(src_type, type, GGML_TYPE_I64, { 31, 3, b, 1 }, { 2, 3, }, 2, v));
test_cases.emplace_back(new test_set_rows(src_type, type, GGML_TYPE_I64, { 33, 5, 1, b }, { 2, 3, }, 1, v));
}
}
}
}
@@ -8013,6 +8017,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
// im2col 2D
test_cases.emplace_back(new test_im2col(GGML_TYPE_F32, GGML_TYPE_F32, GGML_TYPE_F32));
test_cases.emplace_back(new test_im2col(GGML_TYPE_F32, GGML_TYPE_F32, GGML_TYPE_F16));
test_cases.emplace_back(new test_im2col(GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_F32));
test_cases.emplace_back(new test_im2col(GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_F16));
for (int s0 : {1, 3}) {
+4
View File
@@ -152,6 +152,10 @@ int main(int argc, char ** argv) {
init_result = common_init_from_params(params);
ctx = init_result->context();
if (!ctx) {
LOG_ERR("failed to initialize params\n");
return 1;
}
} else {
#ifdef LLAMA_HF_FETCH
auto [hf_repo, hf_quant] = common_download_split_repo_tag(params.model.hf_repo);
+33 -38
View File
@@ -220,8 +220,6 @@ struct server_slot {
return false;
}
GGML_ASSERT(prompt.data.size() == 0);
const size_t cur_size_tgt = llama_state_seq_get_size_ext(ctx_tgt, id, LLAMA_STATE_SEQ_FLAGS_NONE);
const size_t cur_size_dft = ctx_dft ? llama_state_seq_get_size_ext(ctx_dft, id, LLAMA_STATE_SEQ_FLAGS_NONE) : 0;
@@ -252,11 +250,7 @@ struct server_slot {
return res;
}
void prompt_clear(bool allow_processing) {
if (!allow_processing) {
GGML_ASSERT(!is_processing());
}
void prompt_clear() {
SLT_TRC(*this, "clearing prompt with %zu tokens\n", prompt.tokens.size());
common_context_seq_rm(ctx_tgt, id, -1, -1);
@@ -264,7 +258,7 @@ struct server_slot {
common_context_seq_rm(ctx_dft, id, -1, -1);
}
prompt.tokens.clear();
prompt.clear();
}
std::vector<common_adapter_lora_info> lora;
@@ -493,7 +487,7 @@ struct server_slot {
// do not keep context of the child slots - the parent's context is enough
if (task->is_child()) {
prompt_clear(false);
prompt_clear();
}
reset();
@@ -1626,7 +1620,7 @@ private:
ret->prompt_save(*prompt_cache);
if (!ret->prompt_load(*prompt_cache, task.tokens)) {
ret->prompt_clear(false);
ret->prompt_clear();
}
prompt_cache->update();
@@ -1658,7 +1652,7 @@ private:
if (slot.prompt.n_tokens() > 0) {
SRV_WRN("purging slot %d with %zu tokens\n", slot.id, slot.prompt.tokens.size());
slot.prompt_clear(false);
slot.prompt_clear();
res = true;
@@ -1691,7 +1685,7 @@ private:
// if lora has changed, check to see if the cache should be cleared
if (lora_should_clear_cache(slot.lora, task_loras)) {
SLT_TRC(slot, "clearing cache for lora change. %zu loras -> %zu loras\n", slot.lora.size(), task.params.lora.size());
slot.prompt.tokens.clear();
slot.prompt.clear();
} else {
SLT_TRC(slot, "keeping cache for alora. %zu target loras\n", task_loras.size());
}
@@ -2405,7 +2399,7 @@ private:
if (params_base.kv_unified) {
// [TAG_IDLE_SLOT_CLEAR]
slot.prompt_clear(false);
slot.prompt_clear();
}
}
}
@@ -2573,12 +2567,12 @@ private:
size_t token_count = 0;
size_t nread = llama_state_seq_load_file(ctx_tgt, filepath.c_str(), slot->id, tokens.data(), tokens.size(), &token_count);
if (nread == 0) {
slot->prompt.tokens.clear(); // KV may already been invalidated?
slot->prompt.clear(); // KV may already been invalidated?
send_error(task, "Unable to restore slot, no available space in KV cache or invalid slot save file", ERROR_TYPE_INVALID_REQUEST);
break;
}
tokens.resize(token_count);
slot->prompt.tokens.clear();
slot->prompt.clear();
slot->prompt.tokens.insert(tokens);
const int64_t t_end = ggml_time_us();
@@ -2615,7 +2609,7 @@ private:
// Erase token cache
const size_t n_erased = slot->prompt.tokens.size();
slot->prompt_clear(false);
slot->prompt_clear();
auto res = std::make_unique<server_task_result_slot_erase>();
res->id = task.id;
@@ -2775,6 +2769,27 @@ private:
abort_all_slots("pre_decode() failed: " + std::string(e.what()));
}
GGML_ASSERT(batch.slot_batched || batch.size() == 0);
if (batch.slot_batched) {
auto & slot_batched = batch.slot_batched;
auto & alora_scale = batch.alora_scale;
auto & alora_disabled_id = batch.alora_disabled_id;
// TODO @ngxson : alora handling is too messy, need to refactor it to be more clear and maintainable
// apply lora, only need to do it once per batch
common_set_adapter_lora(ctx_tgt, slot_batched->lora);
// if the lora is temporarily disabled for an alora, re-enable it
// for next time
if (alora_scale > 0.0f) {
SRV_DBG("re-enabling alora with scale %f\n", alora_scale);
slot_batched->lora[alora_disabled_id].scale = alora_scale;
}
llama_set_embeddings(ctx_tgt, slot_batched->need_embd());
}
llama_batch batch_view;
int32_t off_next = 0;
int32_t n_batch = llama_n_batch(ctx_tgt);
@@ -2814,7 +2829,6 @@ private:
abort_all_slots("post_decode() failed: " + std::string(e.what()));
break; // stop any further processing
}
}
}
@@ -2880,7 +2894,7 @@ private:
new_tokens.resize(slot.prompt.tokens.size() - n_discard);
slot.prompt.tokens.clear();
slot.prompt.clear();
slot.prompt.tokens.insert(new_tokens);
}
@@ -3556,25 +3570,6 @@ private:
bool decode(int32_t & n_batch, int32_t off, llama_batch & batch_view) {
SRV_DBG("n_batch (effective) = %d, off = %d\n", n_batch, off);
auto & slot_batched = batch.slot_batched;
auto & alora_scale = batch.alora_scale;
auto & alora_disabled_id = batch.alora_disabled_id;
// TODO @ngxson : alora handling is too messy, need to refactor it to be more clear and maintainable
if (slot_batched) {
// apply lora, only need to do it once per batch
common_set_adapter_lora(ctx_tgt, slot_batched->lora);
// if the lora is temporarily disabled for an alora, re-enable it
// for next time
if (alora_scale > 0.0f) {
SRV_DBG("re-enabling alora with scale %f\n", alora_scale);
slot_batched->lora[alora_disabled_id].scale = alora_scale;
}
llama_set_embeddings(ctx_tgt, slot_batched->need_embd());
}
if (batch.size() == 0) {
SRV_WRN("%s", "no tokens to decode\n");
@@ -3622,7 +3617,7 @@ private:
// note: it's complicated to keep track of how much of the current batch has been
// processed before the error occurred, so we simply clear the entire context
slot.prompt_clear(false);
slot.prompt_clear();
}
}
+28 -5
View File
@@ -47,6 +47,16 @@ static void log_server_request(const httplib::Request & req, const httplib::Resp
SRV_DBG("response: %s\n", res.body.c_str());
}
// returns true if the Origin header value's host is localhost / 127.0.0.1 / ::1 (any port)
static bool origin_is_localhost(const std::string & origin) {
try {
const std::string host = common_http_parse_url(origin).host;
return host == "localhost" || host == "127.0.0.1" || host == "::1";
} catch (const std::exception &) {
return false;
}
}
// For Google Cloud Platform deployment compatibility
struct gcp_params {
bool enabled;
@@ -266,13 +276,26 @@ bool server_http_context::init(const common_params & params) {
};
// register server middlewares
srv->set_pre_routing_handler([middleware_validate_api_key, middleware_server_state](const httplib::Request & req, httplib::Response & res) {
res.set_header("Access-Control-Allow-Origin", req.get_header_value("Origin"));
srv->set_pre_routing_handler([&params, middleware_validate_api_key, middleware_server_state](const httplib::Request & req, httplib::Response & res) {
if (params.cors_credentials && params.cors_origins == "*") {
// special case: echo back the Origin header to allow any origin to access the server with credentials
res.set_header("Access-Control-Allow-Origin", req.get_header_value("Origin"));
} else if (params.cors_origins == "localhost") {
// special case: only reflect the Origin header if it is a localhost origin
std::string origin = req.get_header_value("Origin");
if (origin_is_localhost(origin)) {
res.set_header("Access-Control-Allow-Origin", origin);
} else {
SRV_WRN("(CORS) skip non-localhost origin: %s\n", origin.c_str());
}
} else {
res.set_header("Access-Control-Allow-Origin", params.cors_origins);
}
// If this is OPTIONS request, skip validation because browsers don't include Authorization header
if (req.method == "OPTIONS") {
res.set_header("Access-Control-Allow-Credentials", "true");
res.set_header("Access-Control-Allow-Methods", "GET, POST");
res.set_header("Access-Control-Allow-Headers", "*");
res.set_header("Access-Control-Allow-Credentials", params.cors_credentials ? "true" : "false");
res.set_header("Access-Control-Allow-Methods", params.cors_methods);
res.set_header("Access-Control-Allow-Headers", params.cors_headers);
res.set_content("", "text/html"); // blank response, no data
return httplib::Server::HandlerResponse::Handled; // skip further processing
}
+14 -12
View File
@@ -1646,16 +1646,16 @@ size_t server_prompt_cache::n_tokens() const {
size_t res = 0;
for (const auto & state : states) {
res += state.n_tokens();
res += state.prompt.n_tokens();
}
return res;
}
server_prompt * server_prompt_cache::alloc(const server_prompt & prompt, size_t state_size_tgt, size_t state_size_dft) {
server_prompt_cache_state * server_prompt_cache::alloc(const server_prompt & prompt, size_t state_size_tgt, size_t state_size_dft) {
// first check if the current state is contained fully in the cache
for (auto it = states.begin(); it != states.end(); ++it) {
const int cur_lcp_len = it->tokens.get_common_prefix(prompt.tokens);
const int cur_lcp_len = it->prompt.tokens.get_common_prefix(prompt.tokens);
if (cur_lcp_len == (int) prompt.tokens.size()) {
SRV_TRC("%s", " - prompt is already in the cache, skipping\n");
@@ -1680,9 +1680,9 @@ server_prompt * server_prompt_cache::alloc(const server_prompt & prompt, size_t
// remove any cached prompts that are fully contained in the current prompt
for (auto it = states.begin(); it != states.end();) {
const int len = it->tokens.get_common_prefix(prompt.tokens);
const int len = it->prompt.tokens.get_common_prefix(prompt.tokens);
if (len == (int) it->tokens.size()) {
if (len == (int) it->prompt.tokens.size()) {
SRV_TRC(" - removing obsolete cached prompt with length %d\n", len);
it = states.erase(it);
@@ -1721,12 +1721,14 @@ server_prompt * server_prompt_cache::alloc(const server_prompt & prompt, size_t
}
states.push_back({
/*.tokens =*/ prompt.tokens.clone(),
/*.data =*/ {
/*.prompt =*/ {
/*.tokens =*/ prompt.tokens.clone(),
/*.checkpoints =*/ prompt.checkpoints,
},
/*.data =*/ {
/*.main =*/ std::move(state_data_tgt),
/*.drft =*/ std::move(state_data_dft),
},
/*.checkpoints =*/ prompt.checkpoints,
});
return &states.back();
@@ -1744,9 +1746,9 @@ bool server_prompt_cache::load(server_prompt & prompt, const server_tokens & tok
// find the most similar cached prompt, that would also preserve the most context
for (auto it = states.begin(); it != states.end(); ++it) {
const int lcp_cur = it->tokens.get_common_prefix(tokens_new);
const int lcp_cur = it->prompt.tokens.get_common_prefix(tokens_new);
const float f_keep_cur = float(lcp_cur) / it->tokens.size();
const float f_keep_cur = float(lcp_cur) / it->prompt.tokens.size();
const float sim_cur = float(lcp_cur) / tokens_new.size();
// don't trash large prompts
@@ -1799,7 +1801,7 @@ bool server_prompt_cache::load(server_prompt & prompt, const server_tokens & tok
}
}
prompt = std::move(*it_best);
prompt = std::move(it_best->prompt);
states.erase(it_best);
}
@@ -1836,6 +1838,6 @@ void server_prompt_cache::update() {
for (const auto & state : states) {
SRV_TRC(" - 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));
(const void *)&state, state.prompt.n_tokens(), state.prompt.checkpoints.size(), state.size() / (1024.0 * 1024.0));
}
}
+29 -24
View File
@@ -584,32 +584,14 @@ struct server_task_result_apply_lora : server_task_result {
virtual json to_json() override;
};
struct server_prompt_data {
std::vector<uint8_t> main;
std::vector<uint8_t> drft;
size_t size() const {
return main.size() + drft.size();
}
};
struct server_prompt {
server_tokens tokens;
server_prompt_data data;
std::list<common_prompt_checkpoint> checkpoints;
size_t size() const {
size_t res = 0;
res += data.size();
for (const auto & ckpt : checkpoints) {
res += ckpt.size();
}
return res;
void clear() {
tokens.clear();
checkpoints.clear();
}
int n_tokens() const {
@@ -619,19 +601,42 @@ struct server_prompt {
server_prompt clone() const {
return server_prompt {
tokens.clone(),
data,
checkpoints,
};
}
};
struct server_prompt_data {
std::vector<uint8_t> main;
std::vector<uint8_t> drft;
size_t size() const {
return main.size() + drft.size();
}
};
struct server_prompt_cache_state {
server_prompt prompt;
server_prompt_data data;
size_t size() const {
size_t res = data.size();
for (const auto & ckpt : prompt.checkpoints) {
res += ckpt.size();
}
return res;
}
};
struct server_prompt_cache {
server_prompt_cache(int32_t limit_size_mib, size_t limit_tokens) {
this->limit_size = 1024ull*1024ull*(limit_size_mib < 0 ? 0 : limit_size_mib);
this->limit_tokens = limit_tokens;
}
std::list<server_prompt> states;
std::list<server_prompt_cache_state> states;
// in bytes, 0 = no limit
size_t limit_size = 0;
@@ -643,7 +648,7 @@ struct server_prompt_cache {
size_t n_tokens() const;
server_prompt * alloc(const server_prompt & prompt, size_t state_size_main, size_t state_size_drft);
server_prompt_cache_state * alloc(const server_prompt & prompt, size_t state_size_main, size_t state_size_drft);
bool load(server_prompt & prompt, const server_tokens & tokens_new, llama_context * ctx_main, llama_context * ctx_drft, int32_t id_slot);
+25 -11
View File
@@ -303,14 +303,24 @@ int llama_server(common_params & params, int argc, char ** argv) {
return res;
};
if (params.cors_origins == "*" && params.api_keys.empty()) {
SRV_WRN("%s", "-----------------\n");
SRV_WRN("%s", "CORS is set to allow all origins ('*') and no API key is set\n");
SRV_WRN("%s", "this can be a security risk (cross-origin attacks)\n");
SRV_WRN("%s", "more info: https://github.com/ggml-org/llama.cpp/pull/25655\n");
SRV_WRN("%s", "-----------------\n");
}
// CORS proxy (EXPERIMENTAL, only used by the Web UI for MCP)
std::vector<std::string> warn_names;
if (is_router_server) {
warn_names.push_back("router mode");
}
if (params.ui_mcp_proxy) {
SRV_WRN("%s", "-----------------\n");
SRV_WRN("%s", "CORS proxy is enabled, do not expose server to untrusted environments\n");
SRV_WRN("%s", "This feature is EXPERIMENTAL and may be removed or changed in future versions\n");
SRV_WRN("%s", "-----------------\n");
ctx_http.get ("/cors-proxy", ex_wrapper(proxy_handler_get));
ctx_http.post("/cors-proxy", ex_wrapper(proxy_handler_post));
warn_names.push_back("MCP proxy (experimental)");
} else {
ctx_http.get ("/cors-proxy", ex_wrapper(res_403));
ctx_http.post("/cors-proxy", ex_wrapper(res_403));
@@ -324,17 +334,24 @@ int llama_server(common_params & params, int argc, char ** argv) {
SRV_ERR("tools setup failed: %s\n", e.what());
return 1;
}
SRV_WRN("%s", "-----------------\n");
SRV_WRN("%s", "Built-in tools are enabled, do not expose server to untrusted environments\n");
SRV_WRN("%s", "This feature is EXPERIMENTAL and may be changed in the future\n");
SRV_WRN("%s", "-----------------\n");
ctx_http.get ("/tools", ex_wrapper(tools.handle_get));
ctx_http.post("/tools", ex_wrapper(tools.handle_post));
warn_names.push_back("built-in tools (experimental)");
} else {
ctx_http.get ("/tools", ex_wrapper(res_403));
ctx_http.post("/tools", ex_wrapper(res_403));
}
if (warn_names.size() > 0) {
SRV_WRN("%s", "-----------------\n");
SRV_WRN("%s", "the following feature(s) are enabled:\n");
for (const auto & name : warn_names) {
SRV_WRN(" %s\n", name.c_str());
}
SRV_WRN("%s", "do not expose the server to untrusted environments\n");
SRV_WRN("%s", "-----------------\n");
}
//
// Handle downloading model
//
@@ -452,9 +469,6 @@ int llama_server(common_params & params, int argc, char ** argv) {
SRV_INF("listening on %s\n", ctx_http.listening_address.c_str());
if (is_router_server) {
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 (!params.models_preset_hf.empty()) {
SRV_WRN( "NOTE: using preset.ini from HF repo '%s'\n", params.models_preset_hf.c_str());
SRV_WRN("%s", " please only use presets that you can trust! Unknown presets may be unsafe\n");
+65 -1
View File
@@ -91,7 +91,7 @@ def test_openai_library_correct_api_key():
("localhost", "Access-Control-Allow-Origin", "localhost"),
("web.mydomain.fr", "Access-Control-Allow-Origin", "web.mydomain.fr"),
("origin", "Access-Control-Allow-Credentials", "true"),
("web.mydomain.fr", "Access-Control-Allow-Methods", "GET, POST"),
("web.mydomain.fr", "Access-Control-Allow-Methods", "GET, POST, DELETE, OPTIONS"),
("web.mydomain.fr", "Access-Control-Allow-Headers", "*"),
])
def test_cors_options(origin: str, cors_header: str, cors_header_value: str):
@@ -107,6 +107,70 @@ def test_cors_options(origin: str, cors_header: str, cors_header_value: str):
assert res.headers[cors_header] == cors_header_value
@pytest.mark.parametrize("origin", [
"http://localhost",
"http://localhost:8080",
"http://127.0.0.1",
"http://127.0.0.1:3000",
"http://[::1]",
"http://[::1]:3000",
])
def test_cors_origins_localhost_reflects(origin: str):
global server
server = ServerPreset.router()
server.cors_origins = "localhost"
server.start()
res = server.make_request("OPTIONS", "/completions", headers={
"Origin": origin,
"Access-Control-Request-Method": "POST",
"Access-Control-Request-Headers": "Authorization",
})
assert res.status_code == 200
assert res.headers["Access-Control-Allow-Origin"] == origin
@pytest.mark.parametrize("origin", [
"http://web.mydomain.fr",
"http://evil.com",
"http://notlocalhost",
"http://localhost.evil.com",
])
def test_cors_origins_localhost_rejects(origin: str):
global server
server = ServerPreset.router()
server.cors_origins = "localhost"
server.start()
res = server.make_request("OPTIONS", "/completions", headers={
"Origin": origin,
"Access-Control-Request-Method": "POST",
"Access-Control-Request-Headers": "Authorization",
})
assert res.status_code == 200
assert "Access-Control-Allow-Origin" not in res.headers
def test_cors_origins_defaults_to_localhost_with_tools_enabled():
global server
server = ServerPreset.router()
server.server_tools = "all"
server.start()
res = server.make_request("OPTIONS", "/completions", headers={
"Origin": "http://localhost:8080",
"Access-Control-Request-Method": "POST",
"Access-Control-Request-Headers": "Authorization",
})
assert res.status_code == 200
assert res.headers["Access-Control-Allow-Origin"] == "http://localhost:8080"
res = server.make_request("OPTIONS", "/completions", headers={
"Origin": "http://evil.com",
"Access-Control-Request-Method": "POST",
"Access-Control-Request-Headers": "Authorization",
})
assert res.status_code == 200
assert "Access-Control-Allow-Origin" not in res.headers
def test_cors_proxy_only_forwards_explicit_proxy_headers():
class CaptureHeadersHandler(BaseHTTPRequestHandler):
def do_GET(self):
+4 -1
View File
@@ -114,6 +114,7 @@ class ServerProcess:
backend_sampling: bool = False
gcp_compat: bool = False
server_tools: str | None = None
cors_origins: str | None = None
# session variables
process: subprocess.Popen | None = None
@@ -170,6 +171,8 @@ class ServerProcess:
server_args.extend(["--models-max", self.models_max])
if self.models_preset:
server_args.extend(["--models-preset", self.models_preset])
if self.cors_origins:
server_args.extend(["--cors-origins", self.cors_origins])
if self.n_batch:
server_args.extend(["--batch-size", self.n_batch])
if self.n_ubatch:
@@ -359,7 +362,7 @@ class ServerProcess:
if parse_body:
try:
result.body = response.json()
except JSONDecodeError:
except (JSONDecodeError, requests.exceptions.JSONDecodeError):
result.body = response.text
else:
result.body = None
+69 -258
View File
@@ -1,5 +1,6 @@
#include "arg.h"
#include "common.h"
//#include "log.h" // TODO: start using log.h
#include "log.h"
#include "llama.h"
#include <clocale>
@@ -8,115 +9,22 @@
#include <fstream>
#include <string>
#include <vector>
#include <iostream> // TODO: remove me
#include <iostream>
#include <sstream>
#if defined(_WIN32)
#define WIN32_LEAN_AND_MEAN
#include <windows.h>
#include <shellapi.h> // For CommandLineToArgvW
#endif
static void print_usage_information(const char * argv0) {
printf("usage: %s [options]\n\n", argv0);
printf("The tokenize program tokenizes a prompt using a given model,\n");
printf("and prints the resulting tokens to standard output.\n\n");
printf("It needs a model file, a prompt, and optionally other flags\n");
printf("to control the behavior of the tokenizer.\n\n");
printf(" The possible options are:\n");
printf("\n");
printf(" -h, --help print this help and exit\n");
printf(" -m MODEL_PATH, --model MODEL_PATH path to model.\n");
printf(" --ids if given, only print numerical token IDs, and not token strings.\n");
printf(" The output format looks like [1, 2, 3], i.e. parseable by Python.\n");
printf(" -f PROMPT_FNAME, --file PROMPT_FNAME read prompt from a file.\n");
printf(" -p PROMPT, --prompt PROMPT read prompt from the argument.\n");
printf(" --stdin read prompt from standard input.\n");
printf(" --no-bos do not ever add a BOS token to the prompt, even if normally the model uses a BOS token.\n");
printf(" --no-escape do not escape input (such as \\n, \\t, etc.).\n");
printf(" --no-parse-special do not parse control tokens.\n");
printf(" --log-disable disable logs. Makes stderr quiet when loading the model.\n");
printf(" --show-count print the total number of tokens.\n");
}
static void print_usage(int argc, char ** argv) {
(void) argc;
static void llama_log_callback_null(ggml_log_level level, const char * text, void * user_data) {
(void) level;
(void) text;
(void) user_data;
}
static std::string read_prompt_from_file(const char * filepath, bool & success) {
success = false;
std::ifstream in(filepath, std::ios::binary);
if (!in) {
fprintf(stderr, "%s: could not open file '%s' for reading: %s\n", __func__, filepath, strerror(errno));
return std::string();
}
// do not assume the file is seekable (e.g. /dev/stdin)
std::stringstream buffer;
buffer << in.rdbuf();
if (in.fail()) {
fprintf(stderr, "%s: could not read the entire file '%s': %s\n", __func__, filepath, strerror(errno));
return std::string();
}
success = true;
return buffer.str();
}
//
// Function: ingest_args(...) -> vector<string>
//
// Takes argc and argv arguments, and converts them to a vector of UTF-8 encoded
// strings, as an STL vector<string>.
//
// In particular, it handles character encoding shenanigans on Windows.
//
// Note: raw_argc and raw_argv are not actually read at all on Windows.
// On Windows we call GetCommandLineW to get the arguments in wchar_t
// format, ignoring the regular argc/argv arguments to main().
//
// TODO: potential opportunity to roll common stuff into common/console.cpp
// in relation to Windows wchar_t shenanigans.
static std::vector<std::string> ingest_args(int raw_argc, char ** raw_argv) {
std::vector<std::string> argv;
// Handle Windows, if given non-ASCII arguments.
// We convert wchar_t arguments into UTF-8 char* on this platform.
// Lets you invoke 'tokenize' on Windows cmd.exe with non-ASCII characters
// without throwing tantrums.
#if defined(_WIN32)
int argc;
const LPWSTR cmdline_wargv = GetCommandLineW();
LPWSTR * wargv = CommandLineToArgvW(cmdline_wargv, &argc);
// silence unused arg warnings
(void) raw_argc;
(void) raw_argv;
for (int i = 0; i < argc; ++i) {
int length_needed = WideCharToMultiByte(CP_UTF8, 0, wargv[i], wcslen(wargv[i]), 0, 0, NULL, NULL);
char * output_buf = (char *) calloc(length_needed+1, sizeof(char));
GGML_ASSERT(output_buf);
WideCharToMultiByte(CP_UTF8, 0, wargv[i], wcslen(wargv[i]), output_buf, length_needed, NULL, NULL);
output_buf[length_needed] = '\0';
argv.push_back(output_buf);
free(output_buf);
}
LocalFree((HLOCAL) wargv);
#else
int argc = raw_argc;
for (int i = 0; i < argc; ++i) {
argv.push_back(raw_argv[i]);
}
#endif
GGML_ASSERT((unsigned int) argc == argv.size());
return argv;
LOG("\nexample usage:\n");
LOG("\n %s -m your_model.gguf -p \"Hello world\"\n", argv[0]);
LOG("\n %s -m your_model.gguf -f prompt.txt --ids\n", argv[0]);
LOG("\n cat prompt.txt | %s -m your_model.gguf --stdin --show-count\n", argv[0]);
LOG("\n");
}
//
@@ -184,166 +92,69 @@ static void write_utf8_cstr_to_stdout(const char * str, bool & invalid_utf8) {
#endif
}
int main(int raw_argc, char ** raw_argv) {
int main(int argc, char ** argv) {
std::setlocale(LC_NUMERIC, "C");
const std::vector<std::string> argv = ingest_args(raw_argc, raw_argv);
const int argc = argv.size();
common_params params;
if (argc <= 1) {
print_usage_information(argv[0].c_str());
common_init();
if (!common_params_parse(argc, argv, params, LLAMA_EXAMPLE_TOKENIZE, print_usage)) {
return 1;
}
//////
// Read out all the command line arguments.
//////
// which prompt source was requested?
// -p/--prompt and -f/--file both end up in params.prompt (common's -f also
// strips a single trailing newline), but -f additionally records the path
// in params.prompt_file, so we use that to tell them apart.
const bool use_stdin = params.tokenize_stdin;
const bool use_file = !params.prompt_file.empty();
// variables where to put any arguments we see.
bool printing_ids = false;
bool no_bos = false;
bool no_escape = false;
bool no_parse_special = false;
bool disable_logging = false;
bool show_token_count = false;
const char * model_path = NULL;
const char * prompt_path = NULL;
const char * prompt_arg = NULL;
// track which arguments were explicitly given
// used for sanity checking down the line
bool model_path_set = false;
bool prompt_path_set = false;
bool prompt_set = false;
bool stdin_set = false;
int iarg = 1;
for (; iarg < argc; ++iarg) {
std::string arg{argv[iarg]};
if (arg == "-h" || arg == "--help") {
print_usage_information(argv[0].c_str());
return 0;
}
else if (arg == "--ids") {
printing_ids = true;
}
else if (arg == "-m" || arg == "--model") {
if (model_path_set) {
fprintf(stderr, "Error: -m or --model specified multiple times.\n");
return 1;
}
model_path = argv[++iarg].c_str();
model_path_set = true;
}
else if (arg == "--no-bos") {
no_bos = true;
}
else if (arg == "--no-escape") {
no_escape = true;
}
else if (arg == "--no-parse-special") {
no_parse_special = true;
}
else if (arg == "-p" || arg == "--prompt") {
if (prompt_set) {
fprintf(stderr, "Error: -p or --prompt specified multiple times.\n");
return 1;
}
prompt_arg = argv[++iarg].c_str();
prompt_set = true;
}
else if (arg == "-f" || arg == "--file") {
if (prompt_path_set) {
fprintf(stderr, "Error: -f or --file specified multiple times.\n");
return 1;
}
prompt_path = argv[++iarg].c_str();
prompt_path_set = true;
}
else if (arg == "--stdin") {
stdin_set = true;
}
else if (arg == "--log-disable") {
disable_logging = true;
}
else if (arg == "--show-count") {
show_token_count = true;
}
else {
fprintf(stderr, "Error: unknown option '%s'\n", argv[iarg].c_str());
return 1;
}
}
//////
// Sanity check the command line arguments.
//////
// Check that we have the required stuff set.
if (model_path_set && model_path == NULL) {
fprintf(stderr, "Error: --model requires an argument.\n");
return 1;
}
if (!model_path_set) {
fprintf(stderr, "Error: must specify --model.\n");
return 1;
}
if (prompt_path_set && prompt_path == NULL) {
fprintf(stderr, "Error: --file requires an argument.\n");
return 1;
}
if (prompt_set && prompt_arg == NULL) {
fprintf(stderr, "Error: --prompt requires an argument.\n");
return 1;
}
const int prompts_set = !!(prompt_path_set) + !!(prompt_set) + !!(stdin_set);
if (prompts_set > 1) {
fprintf(stderr, "Error: --stdin, --file and --prompt are mutually exclusive.\n");
return 1;
}
// Must have some prompt.
if (prompts_set == 0) {
fprintf(stderr, "Error: must specify one of: --stdin, --file or --prompt.\n");
// sanity check: --stdin is mutually exclusive with -f/--file and -p/--prompt
if (use_stdin && (use_file || !params.prompt.empty())) {
LOG_ERR("error: --stdin is mutually exclusive with --file and --prompt\n");
return 1;
}
GGML_ASSERT(model_path);
GGML_ASSERT(prompt_path || prompt_arg || stdin_set);
//////
// Figure out where will the prompt come from.
//////
// must have some prompt
if (!use_stdin && !use_file && params.prompt.empty()) {
LOG_ERR("error: must specify one of: --stdin, --file or --prompt\n");
return 1;
}
std::string prompt;
if (prompt_path_set) {
bool success = false;
prompt = read_prompt_from_file(prompt_path, success);
if (!success) {
if (use_file) {
// read the file verbatim: common's -f handler strips a single trailing
// newline, but for a tokenizer the input bytes must be preserved exactly
// (a trailing newline is itself a token). escapes are applied locally
// to match the behavior of -p/--prompt and --stdin.
std::ifstream in(params.prompt_file, std::ios::binary);
if (!in) {
LOG_ERR("error: could not open file '%s' for reading\n", params.prompt_file.c_str());
return 1;
}
} else if (prompt_set) {
prompt = prompt_arg;
} else {
GGML_ASSERT(stdin_set);
// we read stdin *after* loading model (early exit if model cannot
// be loaded, which can be a nicer user experience)
}
//////
// Start actually doing the tokenizing stuff.
//////
if (disable_logging) {
llama_log_set(llama_log_callback_null, NULL);
std::stringstream ss;
ss << in.rdbuf();
prompt = ss.str();
if (params.escape) {
string_process_escapes(prompt);
}
} else if (!use_stdin) {
// -p/--prompt is already escape-processed by common_params_parse()
// (controlled by --escape/--no-escape), so use it verbatim here.
prompt = params.prompt;
}
// else: we read stdin *after* loading the model (early exit if the
// model cannot be loaded, which is a nicer user experience)
llama_backend_init();
// load only the vocabulary (no weights), since tokenizing does not need them
llama_model_params model_params = llama_model_default_params();
model_params.vocab_only = true;
llama_model * model = llama_model_load_from_file(model_path, model_params);
llama_model * model = llama_model_load_from_file(params.model.path.c_str(), model_params);
if (!model) {
fprintf(stderr, "Error: could not load model from file '%s'.\n", model_path);
LOG_ERR("error: could not load model from file '%s'.\n", params.model.path.c_str());
return 1;
}
@@ -352,42 +163,41 @@ int main(int raw_argc, char ** raw_argv) {
llama_context_params ctx_params = llama_context_default_params();
llama_context * ctx = llama_init_from_model(model, ctx_params);
if (!ctx) {
fprintf(stderr, "Error: could not create context.\n");
LOG_ERR("error: could not create context.\n");
return 1;
}
// read entire prompt from stdin?
if (stdin_set) {
GGML_ASSERT(!prompt_path_set && !prompt_set);
if (params.tokenize_stdin) {
std::stringstream stdin_buffer;
stdin_buffer << std::cin.rdbuf();
if (std::cin.fail()) {
fprintf(stderr, "Error: could not read the entire standard input.\n");
LOG_ERR("error: could not read the entire standard input.\n");
return 1;
}
prompt = stdin_buffer.str();
// stdin is not seen by common_params_parse(), so apply escape handling
// here to match the behavior of -p/--prompt and -f/--file.
if (params.escape) {
string_process_escapes(prompt);
}
}
const bool model_wants_add_bos = llama_vocab_get_add_bos(vocab);
const bool add_bos = model_wants_add_bos && !no_bos;
const bool parse_special = !no_parse_special;
const bool escape = !no_escape;
if (escape) {
string_process_escapes(prompt);
}
const bool add_bos = model_wants_add_bos && !params.tokenize_no_bos;
const bool parse_special = params.parse_special;
std::vector<llama_token> tokens;
tokens = common_tokenize(vocab, prompt, add_bos, parse_special);
if (printing_ids) {
if (params.tokenize_ids) {
printf("[");
}
for (int i = 0; i < (int) tokens.size(); i++) {
if (printing_ids) {
if (params.tokenize_ids) {
if (i > 0) {
printf(", ");
}
@@ -404,13 +214,14 @@ int main(int raw_argc, char ** raw_argv) {
}
}
if (printing_ids) {
if (params.tokenize_ids) {
printf("]\n");
}
if (show_token_count) {
if (params.tokenize_show_count) {
printf("Total number of tokens: %zu\n", tokens.size());
}
// silence valgrind
llama_free(ctx);
llama_model_free(model);
@@ -17,10 +17,10 @@
let { onMcpSettingsClick }: Props = $props();
let mcpSearchQuery = $state('');
let allMcpServers = $derived(mcpStore.getServers());
let mcpServers = $derived(mcpStore.visibleMcpServers);
// Every configured server is listed; `enabled` is an on/off state,
// not a visibility filter, so a disabled server stays toggleable.
let mcpServers = $derived(mcpStore.getServers());
let hasMcpServers = $derived(mcpServers.length > 0);
// let hasAnyMcpServers = $derived(allMcpServers.length > 0);
let filteredMcpServers = $derived.by(() => {
const query = mcpSearchQuery.toLowerCase().trim();
if (!query) return mcpServers;
@@ -46,7 +46,7 @@
function handleMcpSubMenuOpen(open: boolean) {
if (open) {
mcpSearchQuery = '';
mcpStore.runHealthChecksForServers(allMcpServers);
mcpStore.runHealthChecksForServers(mcpServers);
}
}
@@ -2,85 +2,31 @@
import { Lightbulb, LightbulbOff, Check, Info } from '@lucide/svelte';
import * as DropdownMenu from '$lib/components/ui/dropdown-menu';
import * as Tooltip from '$lib/components/ui/tooltip';
import { ReasoningEffort } from '$lib/enums';
import { REASONING_EFFORT_TOKENS } from '$lib/constants/reasoning-effort-tokens';
import { REASONING_EFFORT_LEVELS } from '$lib/constants/reasoning-effort';
import type { ReasoningEffortLevel } from '$lib/types';
import {
modelsStore,
checkModelSupportsThinking,
supportsThinking,
propsCacheVersion,
loadedModelIds
} from '$lib/stores/models.svelte';
import { chatStore } from '$lib/stores/chat.svelte';
import { conversationsStore, activeMessages } from '$lib/stores/conversations.svelte';
import { isRouterMode } from '$lib/stores/server.svelte';
import type { DatabaseMessage } from '$lib/types/database';
import { useReasoningMenu } from '$lib/hooks/use-reasoning-menu.svelte';
let subOpen = $state(false);
let conversationModel = $derived(
chatStore.getConversationModel(activeMessages() as DatabaseMessage[])
);
let modelSupportsThinkingFromMessages = $derived.by(() => {
const modelId = isRouterMode() ? modelsStore.selectedModelName || conversationModel : null;
if (!modelId) return false;
const messages = conversationsStore.activeMessages;
return messages.some(
(m) => m.role === 'assistant' && m.model === modelId && !!m.reasoningContent
);
});
let modelSupportsThinking = $derived.by(() => {
loadedModelIds();
propsCacheVersion();
if (isRouterMode()) {
const modelId = modelsStore.selectedModelName || conversationModel;
return checkModelSupportsThinking(modelId ?? '') || modelSupportsThinkingFromMessages;
}
return supportsThinking() || modelSupportsThinkingFromMessages;
});
let thinkingEnabled = $derived(conversationsStore.getThinkingEnabled());
let currentEffort = $derived(conversationsStore.getReasoningEffort());
let isOff = $derived(!thinkingEnabled);
function isSelected(item: ReasoningEffortLevel): boolean {
if (item.isOff) return isOff;
return thinkingEnabled && currentEffort === item.value;
}
function handleSelection(item: ReasoningEffortLevel) {
if (item.isOff) {
conversationsStore.setThinkingEnabled(false);
} else {
conversationsStore.setThinkingEnabled(true);
conversationsStore.setReasoningEffort(item.value as ReasoningEffort);
}
subOpen = false;
}
const reasoning = useReasoningMenu();
</script>
{#if modelSupportsThinking}
{#if reasoning.modelSupportsThinking}
<DropdownMenu.Sub bind:open={subOpen}>
<DropdownMenu.SubTrigger class="flex cursor-pointer items-center gap-2">
{#if thinkingEnabled}
{#if reasoning.thinkingEnabled}
<Lightbulb class="h-4 w-4 shrink-0 text-amber-400" />
{:else}
<LightbulbOff class="h-4 w-4 shrink-0 text-muted-foreground" />
{/if}
<span class="text-sm inline-flex gap-2 {!thinkingEnabled ? 'text-muted-foreground' : ''}">
<span
class="text-sm inline-flex gap-2 {!reasoning.thinkingEnabled
? 'text-muted-foreground'
: ''}"
>
Reasoning
<span class="capitalize text-muted-foreground">
{thinkingEnabled ? currentEffort : 'off'}
{reasoning.thinkingEnabled ? reasoning.currentEffort : 'off'}
</span>
</span>
</DropdownMenu.SubTrigger>
@@ -88,14 +34,18 @@
<DropdownMenu.SubContent
class="w-60 bg-popover p-1.5 text-popover-foreground shadow-md outline-none"
>
{#each REASONING_EFFORT_LEVELS as level (level.value)}
{#each reasoning.levels as level (level.value)}
{@const tokenLabel = reasoning.tokenLabel(level)}
<button
type="button"
class="flex w-full cursor-pointer items-center gap-3 rounded-md px-2 py-1.75 text-left text-sm transition-colors hover:bg-accent"
class:bg-accent={isSelected(level)}
onclick={() => handleSelection(level)}
class:bg-accent={reasoning.isSelected(level)}
onclick={() => {
reasoning.select(level);
subOpen = false;
}}
>
{#if isSelected(level)}
{#if reasoning.isSelected(level)}
<Check class="h-4 w-4 shrink-0 text-foreground" />
{:else}
<div class="h-4 w-4 shrink-0"></div>
@@ -103,11 +53,9 @@
<span class="flex-1">{level.label}</span>
{#if !level.isOff}
{#if tokenLabel}
<span class="text-[11px] text-muted-foreground opacity-60">
{REASONING_EFFORT_TOKENS[level.value] === -1
? 'Unlimited'
: `Max ${REASONING_EFFORT_TOKENS[level.value].toLocaleString()} tokens`}
{tokenLabel}
</span>
{/if}
@@ -10,10 +10,18 @@
import { ATTACHMENT_FILE_ITEMS } from '$lib/constants/attachment-menu';
import { useAttachmentMenu } from '$lib/hooks/use-attachment-menu.svelte';
import { useToolsPanel } from '$lib/hooks/use-tools-panel.svelte';
import { useReasoningMenu } from '$lib/hooks/use-reasoning-menu.svelte';
import { conversationsStore } from '$lib/stores/conversations.svelte';
import { mcpStore } from '$lib/stores/mcp.svelte';
import { McpLogo } from '$lib/components/app';
import { PencilRuler, ChevronDown, ChevronRight } from '@lucide/svelte';
import {
PencilRuler,
ChevronDown,
ChevronRight,
Lightbulb,
LightbulbOff,
Check
} from '@lucide/svelte';
import { HealthCheckStatus } from '$lib/enums';
import { AttachmentAction } from '$lib/enums/attachment.enums';
@@ -48,6 +56,7 @@
}: Props = $props();
let sheetOpen = $state(false);
let reasoningExpanded = $state(false);
let filesExpanded = $state(true);
let toolsExpanded = $state(false);
let mcpExpanded = $state(false);
@@ -67,6 +76,7 @@
);
const toolsPanel = useToolsPanel();
const reasoning = useReasoningMenu();
const sheetItemClass =
'flex w-full items-center gap-3 rounded-md px-3 py-2.5 text-left text-sm transition-colors hover:bg-accent active:bg-accent disabled:cursor-not-allowed disabled:opacity-50';
@@ -74,7 +84,7 @@
const sheetItemRowClass =
'flex w-full items-center justify-between gap-2 rounded-md px-3 py-2 text-left text-sm transition-colors hover:bg-accent';
let visibleMcpServers = $derived(mcpStore.visibleMcpServers);
let mcpServers = $derived(mcpStore.getServers());
</script>
<div class="flex items-center gap-1 {className}">
@@ -91,6 +101,63 @@
</Sheet.Header>
<div class="flex flex-col gap-1 px-1.5 pb-2">
{#if reasoning.modelSupportsThinking}
<Collapsible.Root
open={reasoningExpanded}
onOpenChange={(open) => (reasoningExpanded = open)}
>
<Collapsible.Trigger class={sheetItemClass}>
{#if reasoningExpanded}
<ChevronDown class="h-4 w-4 shrink-0" />
{:else}
<ChevronRight class="h-4 w-4 shrink-0" />
{/if}
{#if reasoning.thinkingEnabled}
<Lightbulb class="h-4 w-4 shrink-0 text-amber-400" />
{:else}
<LightbulbOff class="h-4 w-4 shrink-0 text-muted-foreground" />
{/if}
<span class="flex-1">Reasoning</span>
<span class="text-xs capitalize text-muted-foreground">
{reasoning.thinkingEnabled ? reasoning.currentEffort : 'off'}
</span>
</Collapsible.Trigger>
<Collapsible.Content>
<div class="flex flex-col gap-0.5 pl-4">
{#each reasoning.levels as level (level.value)}
{@const tokenLabel = reasoning.tokenLabel(level)}
<button
type="button"
class={sheetItemRowClass}
class:bg-accent={reasoning.isSelected(level)}
onclick={() => reasoning.select(level)}
>
<div class="flex min-w-0 items-center gap-3">
{#if reasoning.isSelected(level)}
<Check class="h-4 w-4 shrink-0 text-foreground" />
{:else}
<div class="h-4 w-4 shrink-0"></div>
{/if}
<span class="text-sm">{level.label}</span>
</div>
{#if tokenLabel}
<span class="shrink-0 text-[11px] text-muted-foreground opacity-60">
{tokenLabel}
</span>
{/if}
</button>
{/each}
</div>
</Collapsible.Content>
</Collapsible.Root>
{/if}
<Collapsible.Root open={filesExpanded} onOpenChange={(open) => (filesExpanded = open)}>
<Collapsible.Trigger class={sheetItemClass}>
{#if filesExpanded}
@@ -151,13 +218,13 @@
<span class="flex-1">MCP Servers</span>
<span class="text-xs text-muted-foreground">
{visibleMcpServers.length} server{visibleMcpServers.length !== 1 ? 's' : ''}
{mcpServers.length} server{mcpServers.length !== 1 ? 's' : ''}
</span>
</Collapsible.Trigger>
<Collapsible.Content>
<div class="flex flex-col gap-0.5 pl-4">
{#each visibleMcpServers as server (server.id)}
{#each mcpServers as server (server.id)}
{@const healthState = mcpStore.getHealthCheckState(server.id)}
{@const hasError = healthState.status === HealthCheckStatus.ERROR}
{@const displayName = mcpStore.getServerLabel(server)}
@@ -200,7 +267,7 @@
</button>
{/each}
{#if visibleMcpServers.length === 0}
{#if mcpServers.length === 0}
<div class="px-3 py-2 text-center text-sm text-muted-foreground">
No MCP servers configured
</div>
@@ -21,7 +21,7 @@
<ChatMessageActionCard icon={ShieldQuestion}>
{#snippet message()}
Allow use of <span class="font-semibold">{toolName}</span>{#if serverLabel}
from <span class="font-semibold">{serverLabel}</span>{/if}?
&nbsp;from <span class="font-semibold">{serverLabel}</span>{/if}?
{/snippet}
{#snippet actions()}
@@ -16,6 +16,7 @@
let newServerUrl = $state('');
let newServerHeaders = $state('');
let newServerUseProxy = $state(false);
let newServerUrlError = $derived.by(() => {
if (!newServerUrl.trim()) return 'URL is required';
try {
@@ -35,6 +36,7 @@
if (!value) {
newServerUrl = '';
newServerHeaders = '';
newServerUseProxy = false;
}
open = value;
onOpenChange?.(value);
@@ -49,7 +51,8 @@
id: newServerId,
enabled: true,
url: newServerUrl.trim(),
headers: newServerHeaders.trim() || undefined
headers: newServerHeaders.trim() || undefined,
useProxy: newServerUseProxy
});
conversationsStore.setMcpServerOverride(newServerId, true);
@@ -74,8 +77,10 @@
<McpServerForm
url={newServerUrl}
headers={newServerHeaders}
useProxy={newServerUseProxy}
onUrlChange={(v) => (newServerUrl = v)}
onHeadersChange={(v) => (newServerHeaders = v)}
onUseProxyChange={(v) => (newServerUseProxy = v)}
urlError={newServerUrl ? newServerUrlError : null}
id="new-server"
/>
@@ -32,7 +32,9 @@
let isHealthChecking = $derived(healthState.status === HealthCheckStatus.CONNECTING);
let isConnected = $derived(healthState.status === HealthCheckStatus.SUCCESS);
let isError = $derived(healthState.status === HealthCheckStatus.ERROR);
let showSkeleton = $derived(isIdle || isHealthChecking);
// Disabled servers stay IDLE (no startup health check), so the body
// skeleton only applies while a check is running or expected to run.
let showSkeleton = $derived(isHealthChecking || (isIdle && server.enabled));
let errorMessage = $derived(
healthState.status === HealthCheckStatus.ERROR ? healthState.message : undefined
);
@@ -22,7 +22,9 @@
let { class: className }: Props = $props();
let servers = $derived(mcpStore.visibleMcpServers);
// Every configured server is listed; `enabled` is an on/off state,
// not a visibility filter, so a disabled server stays toggleable.
let servers = $derived(mcpStore.getServers());
let isAddingServer = $state(false);
@@ -58,9 +60,14 @@
// Each card decides for itself whether to render based on its own
// health-check state, so adding a server only flashes the new card
// (not every other already-loaded card) until its health check resolves.
function isServerPending(serverId: string): boolean {
// Disabled servers never receive a startup health check, so IDLE only
// counts as pending when the server is enabled; otherwise the real card
// renders and keeps the enable toggle reachable.
function isServerPending(serverId: string, enabled: boolean): boolean {
const status = mcpStore.getHealthCheckState(serverId).status;
return status === HealthCheckStatus.IDLE || status === HealthCheckStatus.CONNECTING;
return (
status === HealthCheckStatus.CONNECTING || (status === HealthCheckStatus.IDLE && enabled)
);
}
</script>
@@ -109,7 +116,7 @@
style="grid-template-columns: repeat(auto-fill, minmax(min(32rem, calc(100dvw - 2rem)), 1fr));"
>
{#each servers as server (server.id)}
{#if isServerPending(server.id)}
{#if isServerPending(server.id, server.enabled)}
<McpServerCardSkeleton />
{:else}
<McpServerCard
+1 -2
View File
@@ -6,8 +6,7 @@ export const NEWLINE_SEPARATOR = '\n';
export const DEFAULT_AGENTIC_CONFIG: AgenticConfig = {
enabled: true,
maxTurns: 100,
maxToolPreviewLines: 25
maxTurns: 100
} as const;
export const REASONING_TAGS = {
+3 -1
View File
@@ -24,8 +24,10 @@ export const MODEL_CUSTOM_QUANTIZATION_PREFIX_RE = /^UD$/i;
/**
* Matches a parameter-count segment, e.g. `7B`, `1.5b`, `120M`.
* The optional leading `E` covers effective-parameter sizes, e.g. Gemma's
* `E2B`/`E4B` (MatFormer models sized by resident params).
*/
export const MODEL_PARAMS_RE = /^\d+(\.\d+)?[BbMmKkTt]$/;
export const MODEL_PARAMS_RE = /^[Ee]?\d+(\.\d+)?[BbMmKkTt]$/;
/**
* Matches an activated-parameter-count segment, e.g. `A10B`, `a2.4b`.
-1
View File
@@ -8,7 +8,6 @@ export const SETTINGS_SECTION_SLUGS = {
PENALTIES: 'penalties',
AGENTIC: 'agentic',
DEVELOPER: 'developer',
MCP: 'mcp',
TOOLS: 'tools',
IMPORT_EXPORT: 'import-export'
} as const;
@@ -59,7 +59,6 @@ export const SETTINGS_KEYS = {
MCP_SERVERS: 'mcpServers',
MCP_REQUEST_TIMEOUT_SECONDS: 'mcpRequestTimeoutSeconds',
AGENTIC_MAX_TURNS: 'agenticMaxTurns',
AGENTIC_MAX_TOOL_PREVIEW_LINES: 'agenticMaxToolPreviewLines',
SHOW_TOOL_CALL_IN_PROGRESS: 'showToolCallInProgress',
// Performance
PRE_ENCODE_CONVERSATION: 'preEncodeConversation',
@@ -24,7 +24,6 @@ import type {
SettingsSection
} from '$lib/types';
import { CLI_FLAGS, DEFAULT_MCP_CONFIG } from '$lib/constants';
import McpLogo from '$lib/components/app/mcp/McpLogo.svelte';
import { SETTINGS_KEYS } from './settings-keys';
import { ROUTES, SETTINGS_SECTION_SLUGS } from './routes';
import { TITLE_GENERATION } from './title-generation';
@@ -36,7 +35,6 @@ export const SETTINGS_SECTION_TITLES = {
PENALTIES: 'Penalties',
AGENTIC: 'Agentic',
TOOLS: 'Tools',
MCP: 'MCP',
IMPORT_EXPORT: 'Import/Export',
DEVELOPER: 'Developer'
} as const;
@@ -635,15 +633,15 @@ const SETTINGS_REGISTRY: Record<string, SettingsSectionEntry> = {
}
},
{
key: SETTINGS_KEYS.AGENTIC_MAX_TOOL_PREVIEW_LINES,
label: 'Max lines per tool preview',
help: 'Number of lines shown in tool output previews (last N lines). Only these previews and the final LLM response persist after the agentic loop completes.',
defaultValue: 25,
key: SETTINGS_KEYS.MCP_REQUEST_TIMEOUT_SECONDS,
label: 'MCP request timeout (seconds)',
help: 'Timeout for individual MCP tool calls.',
defaultValue: DEFAULT_MCP_CONFIG.requestTimeoutSeconds,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.AGENTIC,
isPositiveInteger: true,
sync: {
serverKey: SETTINGS_KEYS.AGENTIC_MAX_TOOL_PREVIEW_LINES,
serverKey: SETTINGS_KEYS.MCP_REQUEST_TIMEOUT_SECONDS,
paramType: SyncableParameterType.NUMBER
}
}
@@ -735,26 +733,6 @@ const SETTINGS_REGISTRY: Record<string, SettingsSectionEntry> = {
}
}
]
},
[SETTINGS_SECTION_SLUGS.MCP]: {
title: SETTINGS_SECTION_TITLES.MCP,
slug: SETTINGS_SECTION_SLUGS.MCP,
icon: McpLogo,
settings: [
{
key: SETTINGS_KEYS.MCP_REQUEST_TIMEOUT_SECONDS,
label: 'Request timeout (seconds)',
help: 'Default timeout for individual MCP tool calls. Can be overridden per server.',
defaultValue: DEFAULT_MCP_CONFIG.requestTimeoutSeconds,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.MCP,
isPositiveInteger: true,
sync: {
serverKey: SETTINGS_KEYS.MCP_REQUEST_TIMEOUT_SECONDS,
paramType: SyncableParameterType.NUMBER
}
}
]
}
} as const;
@@ -0,0 +1,96 @@
import { ReasoningEffort } from '$lib/enums';
import { REASONING_EFFORT_LEVELS } from '$lib/constants/reasoning-effort';
import { REASONING_EFFORT_TOKENS } from '$lib/constants/reasoning-effort-tokens';
import type { ReasoningEffortLevel } from '$lib/types';
import type { DatabaseMessage } from '$lib/types/database';
import {
modelsStore,
checkModelSupportsThinking,
supportsThinking,
propsCacheVersion,
loadedModelIds
} from '$lib/stores/models.svelte';
import { chatStore } from '$lib/stores/chat.svelte';
import { conversationsStore, activeMessages } from '$lib/stores/conversations.svelte';
import { isRouterMode } from '$lib/stores/server.svelte';
export interface UseReasoningMenuReturn {
readonly modelSupportsThinking: boolean;
readonly thinkingEnabled: boolean;
readonly currentEffort: ReasoningEffort;
readonly levels: ReasoningEffortLevel[];
isSelected(level: ReasoningEffortLevel): boolean;
tokenLabel(level: ReasoningEffortLevel): string | null;
select(level: ReasoningEffortLevel): void;
}
/**
* Shared reactive state and helpers for the reasoning effort menu.
*
* Used by both the desktop dropdown (`ChatFormActionAddReasoningSubmenu`)
* and the mobile sheet (`ChatFormActionAddSheet`) to avoid duplicating the
* thinking-support derivation and the effort selection logic.
*/
export function useReasoningMenu(): UseReasoningMenuReturn {
const conversationModel = $derived(
chatStore.getConversationModel(activeMessages() as DatabaseMessage[])
);
// a router chat can carry reasoning from an earlier turn before the props
// cache is primed, so a model that already produced thinking still qualifies
const modelSupportsThinkingFromMessages = $derived.by(() => {
const modelId = isRouterMode() ? modelsStore.selectedModelName || conversationModel : null;
if (!modelId) return false;
return conversationsStore.activeMessages.some(
(m) => m.role === 'assistant' && m.model === modelId && !!m.reasoningContent
);
});
const modelSupportsThinking = $derived.by(() => {
loadedModelIds();
propsCacheVersion();
if (isRouterMode()) {
const modelId = modelsStore.selectedModelName || conversationModel;
return checkModelSupportsThinking(modelId ?? '') || modelSupportsThinkingFromMessages;
}
return supportsThinking() || modelSupportsThinkingFromMessages;
});
const thinkingEnabled = $derived(conversationsStore.getThinkingEnabled());
const currentEffort = $derived(conversationsStore.getReasoningEffort());
return {
get modelSupportsThinking() {
return modelSupportsThinking;
},
get thinkingEnabled() {
return thinkingEnabled;
},
get currentEffort() {
return currentEffort;
},
get levels() {
return REASONING_EFFORT_LEVELS;
},
isSelected(level: ReasoningEffortLevel): boolean {
if (level.isOff) return !thinkingEnabled;
return thinkingEnabled && currentEffort === level.value;
},
tokenLabel(level: ReasoningEffortLevel): string | null {
if (level.isOff) return null;
const tokens = REASONING_EFFORT_TOKENS[level.value];
return tokens === -1 ? 'Unlimited' : `Max ${tokens.toLocaleString()} tokens`;
},
select(level: ReasoningEffortLevel): void {
if (level.isOff) {
conversationsStore.setThinkingEnabled(false);
return;
}
conversationsStore.setThinkingEnabled(true);
conversationsStore.setReasoningEffort(level.value as ReasoningEffort);
}
};
}
+24 -1
View File
@@ -692,8 +692,31 @@ export class MCPService {
this.createLog(MCPConnectionPhase.INITIALIZING, 'Sending initialize request...')
);
// The SDK timeout only covers the initialize request, not transport.start(),
// which can hang forever on an unreachable host (SSE endpoint wait, WebSocket
// handshake, proxied fetch). This race bounds the whole handshake and closes
// the transport on expiry so the underlying fetch or socket is aborted.
const handshakeTimeoutMs =
serverConfig.handshakeTimeoutMs ?? DEFAULT_MCP_CONFIG.connectionTimeoutMs;
try {
await client.connect(transport);
let handshakeTimer: ReturnType<typeof setTimeout> | undefined;
const handshakeDeadline = new Promise<never>((_, reject) => {
handshakeTimer = setTimeout(() => {
void transport.close().catch(() => {});
reject(new Error(`Connection timed out after ${Math.round(handshakeTimeoutMs / 1000)}s`));
}, handshakeTimeoutMs);
});
try {
await Promise.race([
client.connect(transport, { timeout: handshakeTimeoutMs }),
handshakeDeadline
]);
} finally {
clearTimeout(handshakeTimer);
}
// Transport diagnostics are only for the initial handshake, not long-lived traffic.
stopPhaseLogging();
client.onerror = runtimeErrorHandler;
+1 -4
View File
@@ -280,16 +280,13 @@ class AgenticStore {
getConfig(settings: SettingsConfigType, perChatOverrides?: McpServerOverride[]): AgenticConfig {
const maxTurns = Number(settings.agenticMaxTurns) || DEFAULT_AGENTIC_CONFIG.maxTurns;
const maxToolPreviewLines =
Number(settings.agenticMaxToolPreviewLines) || DEFAULT_AGENTIC_CONFIG.maxToolPreviewLines;
const hasTools =
mcpStore.hasEnabledServers(perChatOverrides) ||
toolsStore.builtinTools.length > 0 ||
toolsStore.customTools.length > 0;
return {
enabled: hasTools && DEFAULT_AGENTIC_CONFIG.enabled,
maxTurns,
maxToolPreviewLines
maxTurns
};
}
+1 -1
View File
@@ -1334,7 +1334,7 @@ class ChatStore {
}
};
const perChatOverrides = conversationsStore.activeConversation?.mcpServerOverrides;
const perChatOverrides = conversationsStore.getAllMcpServerOverrides();
{
const agenticResult = await agenticStore.runAgenticFlow({
+23 -30
View File
@@ -243,9 +243,9 @@ class ConversationsStore {
const conversationName = name || `Chat ${new Date().toLocaleString()}`;
const conversation = await DatabaseService.createConversation(conversationName);
// New conversations inherit per-server enabled defaults directly from
// `mcpServers[i].enabled` (see #checkServerEnabled). No per-conversation
// override list needs to be seeded.
// No MCP override list is seeded: getAllMcpServerOverrides resolves
// servers without a per-conversation override to `mcpServers[i].enabled`,
// and only explicit toggles are stored on the conversation.
// Inherit global thinking/reasoning defaults into the new conversation
const thinkingEnabled = this.getThinkingEnabled();
@@ -601,48 +601,41 @@ class ConversationsStore {
*/
/**
/**
* Resolve the per-server enabled value when no active conversation exists.
* The default for new chats is the server's own `enabled` flag in `mcpServers`.
* Resolve the default enabled value for a server: its own `enabled`
* flag in `mcpServers`, so the global on/off state lives in one place.
*/
#getDefaultOverrideForNoConversation(serverId: string): McpServerOverride | undefined {
#getDefaultOverride(serverId: string): McpServerOverride | undefined {
const server = mcpStore.getServers().find((s) => s.id === serverId);
if (!server) return undefined;
return { serverId, enabled: server.enabled };
}
/**
* Default overrides for new chats are derived from `mcpServers[i].enabled`,
* so the global on/off state lives in one place.
*/
#getAllDefaultOverridesForNoConversation(): McpServerOverride[] {
return mcpStore.getServers().map((s) => ({ serverId: s.id, enabled: s.enabled }));
}
/**
* Gets MCP server override for a specific server in the active conversation.
* Falls back to `mcpServers[i].enabled` if no active conversation exists.
* Gets the effective MCP server override for a specific server.
* A per-conversation override wins when present; a server without one
* resolves to its `mcpServers[i].enabled` default.
* @param serverId - The server ID to check
* @returns The override if set, undefined if no matching server
* @returns The effective override, undefined if no matching server
*/
getMcpServerOverride(serverId: string): McpServerOverride | undefined {
if (this.activeConversation) {
return this.activeConversation.mcpServerOverrides?.find(
(o: McpServerOverride) => o.serverId === serverId
);
}
return this.#getDefaultOverrideForNoConversation(serverId);
const override = this.activeConversation?.mcpServerOverrides?.find(
(o: McpServerOverride) => o.serverId === serverId
);
if (override) return override;
return this.#getDefaultOverride(serverId);
}
/**
* Get all MCP server overrides for the current conversation.
* When no active conversation, derives from `mcpServers[i].enabled`.
* Gets the effective override list for the current conversation:
* one entry per configured server, resolved per server. The stored
* per-conversation list is sparse and only holds explicit toggles.
*/
getAllMcpServerOverrides(): McpServerOverride[] {
if (this.activeConversation?.mcpServerOverrides) {
return this.activeConversation.mcpServerOverrides;
}
return this.#getAllDefaultOverridesForNoConversation();
const overrides = this.activeConversation?.mcpServerOverrides;
return mcpStore.getServers().map((s) => {
const override = overrides?.find((o: McpServerOverride) => o.serverId === s.id);
return { serverId: s.id, enabled: override?.enabled ?? s.enabled };
});
}
/**
+17 -21
View File
@@ -148,15 +148,22 @@ class MCPStore {
enabled: Boolean((entry as { enabled?: unknown })?.enabled),
url,
name: (entry as { name?: string })?.name,
requestTimeoutSeconds:
(entry as { requestTimeoutSeconds?: number })?.requestTimeoutSeconds ??
DEFAULT_MCP_CONFIG.requestTimeoutSeconds,
headers: headers || undefined,
useProxy: Boolean((entry as { useProxy?: unknown })?.useProxy)
} satisfies MCPServerSettingsEntry;
});
}
/**
* Request timeout in milliseconds, read live from the global setting
* so a change in Settings applies to every server immediately.
*/
#requestTimeoutMs(): number {
const seconds =
Number(config().mcpRequestTimeoutSeconds) || DEFAULT_MCP_CONFIG.requestTimeoutSeconds;
return Math.round(seconds * 1000);
}
/**
* Builds server configuration from a settings entry.
*/
@@ -183,7 +190,7 @@ class MCPStore {
url: entry.url,
transport: detectMcpTransportFromUrl(entry.url),
handshakeTimeoutMs: connectionTimeoutMs,
requestTimeoutMs: Math.round(entry.requestTimeoutSeconds * 1000),
requestTimeoutMs: this.#requestTimeoutMs(),
headers,
useProxy: entry.useProxy
};
@@ -191,15 +198,15 @@ class MCPStore {
/**
* Checks if a server is enabled for a given chat.
* Only per-chat overrides (persisted in localStorage for new chats,
* or in IndexedDB for existing conversations) control enabled state.
* A per-chat override wins when present; a server without one resolves
* to its own `enabled` flag in `mcpServers`.
*/
#checkServerEnabled(
server: MCPServerSettingsEntry,
perChatOverrides?: McpServerOverride[]
): boolean {
const override = perChatOverrides?.find((o) => o.serverId === server.id);
return override?.enabled ?? false;
return override?.enabled ?? server.enabled;
}
/**
@@ -230,7 +237,7 @@ class MCPStore {
protocolVersion: DEFAULT_MCP_CONFIG.protocolVersion,
capabilities: DEFAULT_MCP_CONFIG.capabilities,
clientInfo: DEFAULT_MCP_CONFIG.clientInfo,
requestTimeoutMs: Math.round(DEFAULT_MCP_CONFIG.requestTimeoutSeconds * 1000),
requestTimeoutMs: this.#requestTimeoutMs(),
servers
};
}
@@ -500,7 +507,7 @@ class MCPStore {
}
addServer(
serverData: Omit<MCPServerSettingsEntry, 'id' | 'requestTimeoutSeconds'> & { id?: string }
serverData: Omit<MCPServerSettingsEntry, 'id'> & { id?: string }
): MCPServerSettingsEntry {
const servers = this.getServers();
const newServer: MCPServerSettingsEntry = {
@@ -509,8 +516,6 @@ class MCPStore {
url: serverData.url.trim(),
name: serverData.name,
headers: serverData.headers?.trim() || undefined,
requestTimeoutSeconds:
Number(config().mcpRequestTimeoutSeconds) || DEFAULT_MCP_CONFIG.requestTimeoutSeconds,
useProxy: serverData.useProxy
};
settingsStore.updateConfig(SETTINGS_KEYS.MCP_SERVERS, JSON.stringify([...servers, newServer]));
@@ -551,14 +556,6 @@ class MCPStore {
});
}
/**
* MCP servers selectable in chat-add UIs and the settings page,
* in the order they were added to the config.
*/
get visibleMcpServers(): MCPServerSettingsEntry[] {
return this.getServers().filter((server) => server.enabled);
}
async ensureInitialized(perChatOverrides?: McpServerOverride[]): Promise<boolean> {
if (!browser) {
return false;
@@ -1226,7 +1223,6 @@ class MCPStore {
id: string;
enabled: boolean;
url: string;
requestTimeoutSeconds: number;
headers?: string;
}[],
skipIfChecked = true,
@@ -1317,7 +1313,7 @@ class MCPStore {
logs: []
});
const timeoutMs = Math.round(server.requestTimeoutSeconds * 1000);
const timeoutMs = this.#requestTimeoutMs();
const headers = this.parseHeaders(server.headers);
try {
+2 -1
View File
@@ -412,7 +412,8 @@ class ToolsStore {
tools: { name: string; description?: string }[];
}[] {
const result: ReturnType<ToolsStore['getMcpToolsFromHealthChecks']> = [];
for (const server of mcpStore.visibleMcpServers) {
for (const server of mcpStore.getServers()) {
if (!server.enabled) continue;
const health = mcpStore.getHealthCheckState(server.id);
if (health.status === HealthCheckStatus.SUCCESS && health.tools.length > 0) {
result.push({
-1
View File
@@ -15,7 +15,6 @@ import type { DatabaseMessage, DatabaseMessageExtra, McpServerOverride } from '.
export interface AgenticConfig {
enabled: boolean;
maxTurns: number;
maxToolPreviewLines: number;
}
/**
-2
View File
@@ -174,7 +174,6 @@ export interface HealthCheckParams {
id: string;
enabled: boolean;
url: string;
requestTimeoutSeconds: number;
headers?: string;
useProxy?: boolean;
}
@@ -220,7 +219,6 @@ export interface MCPServerDisplayInfo {
export type MCPServerSettingsEntry = MCPServerDisplayInfo & {
enabled: boolean;
requestTimeoutSeconds: number;
headers?: string;
iconUrl?: string;
useProxy?: boolean;
-5
View File
@@ -9,7 +9,6 @@ import {
MimeTypeText
} from '$lib/enums';
import {
DEFAULT_MCP_CONFIG,
MCP_SERVER_ID_PREFIX,
IMAGE_FILE_EXTENSION_REGEX,
CODE_FILE_EXTENSION_REGEX,
@@ -64,7 +63,6 @@ export function detectMcpTransportFromUrl(url: string): MCPTransportType {
/**
* Parses MCP server settings from a JSON string or array.
* Preserves per-server requestTimeoutSeconds if stored, otherwise falls back to the global default.
* @param rawServers - The raw servers to parse
* @returns An empty array if the input is invalid.
*/
@@ -103,9 +101,6 @@ export function parseMcpServerSettings(rawServers: unknown): MCPServerSettingsEn
enabled: Boolean((entry as { enabled?: unknown })?.enabled),
url,
name: (entry as { name?: string })?.name,
requestTimeoutSeconds:
(entry as { requestTimeoutSeconds?: number })?.requestTimeoutSeconds ??
DEFAULT_MCP_CONFIG.requestTimeoutSeconds,
headers: headers || undefined,
useProxy: Boolean((entry as { useProxy?: unknown })?.useProxy)
} satisfies MCPServerSettingsEntry;
@@ -36,6 +36,11 @@ describe('parseModelId', () => {
expect(parseModelId('model-100b:q4_k_m')).toMatchObject({ params: '100B' });
});
it('extracts effective parameters correctly', () => {
expect(parseModelId('model-E4B-BF16')).toMatchObject({ params: 'E4B' });
expect(parseModelId('model-e2b:q4_k_m')).toMatchObject({ params: 'E2B' });
});
it('extracts activated parameters correctly', () => {
expect(parseModelId('model-100B-A10B-BF16')).toMatchObject({ activatedParams: 'A10B' });
expect(parseModelId('model-100B-A10B:Q4_K_M')).toMatchObject({ activatedParams: 'A10B' });
@@ -1,6 +1,6 @@
import { describe, expect, it, vi } from 'vitest';
import { parseMcpServerSettings } from '$lib/utils/mcp';
import { DEFAULT_MCP_CONFIG, MCP_SERVER_ID_PREFIX } from '$lib/constants/mcp';
import { MCP_SERVER_ID_PREFIX } from '$lib/constants/mcp';
/**
* Tests for the mcpServers settings parser.
@@ -58,24 +58,16 @@ describe('parseMcpServerSettings', () => {
expect(parsed[2]?.id).toBe('custom-3');
});
it('falls back to the configured default requestTimeoutSeconds only for nullish values', () => {
const fallback = DEFAULT_MCP_CONFIG.requestTimeoutSeconds;
it('does not emit a per-server timeout, the request timeout is a live global setting', () => {
// A stored per-server requestTimeoutSeconds was never editable in
// any UI and froze the global setting at server creation time,
// making the Settings value a no-op for existing servers. The
// parser drops the field so the global applies live everywhere.
const parsed = parseMcpServerSettings(
JSON.stringify([
{ id: 'a', url: 'https://a.test' },
{ id: 'b', url: 'https://b.test', requestTimeoutSeconds: undefined },
{ id: 'c', url: 'https://c.test', requestTimeoutSeconds: 0 },
{ id: 'd', url: 'https://d.test', requestTimeoutSeconds: 45 }
])
JSON.stringify([{ id: 'a', url: 'https://a.test', requestTimeoutSeconds: 45 }])
);
// The parser uses ?? for timeout fallback, which only triggers on
// null/undefined. Explicit 0 is preserved at face value.
expect(parsed[0]?.requestTimeoutSeconds).toBe(fallback);
expect(parsed[1]?.requestTimeoutSeconds).toBe(fallback);
expect(parsed[2]?.requestTimeoutSeconds).toBe(0);
expect(parsed[3]?.requestTimeoutSeconds).toBe(45);
expect(parsed[0]).not.toHaveProperty('requestTimeoutSeconds');
});
it('treats whitespace-only headers strings as undefined', () => {
@@ -108,6 +100,22 @@ describe('parseMcpServerSettings', () => {
expect(parsed[3]?.useProxy).toBe(true);
});
it('keeps disabled entries in the list, enabled is state and never a visibility filter', () => {
// Regression guard for issue #25625: filtering the server list on
// `enabled` hides a toggled-off server from every UI surface with
// no way to re-enable it. Any list derived from this parser must
// contain disabled entries.
const parsed = parseMcpServerSettings(
JSON.stringify([
{ id: 'on', url: 'https://on.test', enabled: true },
{ id: 'off', url: 'https://off.test', enabled: false }
])
);
expect(parsed.map((entry) => entry.id)).toEqual(['on', 'off']);
expect(parsed[1]?.enabled).toBe(false);
});
it('preserves input order when mapping entries', () => {
const source = [
{ id: 'gamma', url: 'https://c.test' },