From 69bf6437914596fbbc4caf09a7ac16f2acdd1a94 Mon Sep 17 00:00:00 2001 From: Rafail Giavrimis <47496212+grafail@users.noreply.github.com> Date: Sat, 8 Aug 2026 05:40:04 +0100 Subject: [PATCH 01/16] CUDA: fix thread/block count in quantized cpy kernel launches (#26731) * CUDA: fix thread/block count in quantized cpy kernel launches * tests: add uneven block count cpy case --- ggml/src/ggml-cuda/cpy.cu | 44 +++++++++++++++++++------------------- tests/test-backend-ops.cpp | 3 +++ 2 files changed, 25 insertions(+), 22 deletions(-) diff --git a/ggml/src/ggml-cuda/cpy.cu b/ggml/src/ggml-cuda/cpy.cu index eb5eb0eb4..fd7ffc0bc 100644 --- a/ggml/src/ggml-cuda/cpy.cu +++ b/ggml/src/ggml-cuda/cpy.cu @@ -253,9 +253,9 @@ static void ggml_cpy_f32_q8_0_cuda( const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) { GGML_ASSERT(ne % QK8_0 == 0); - const int64_t num_blocks = ne / QK8_0; + const int64_t num_blocks = (ne/QK8_0 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE; GGML_ASSERT(num_blocks <= INT_MAX); - cpy_f32_q<<>> + cpy_f32_q<<>> (cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13); } @@ -264,9 +264,9 @@ static void ggml_cpy_q8_0_f32_cuda( const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) { - const int64_t num_blocks = ne; + const int64_t num_blocks = (ne/QK8_0 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE; GGML_ASSERT(num_blocks <= INT_MAX); - cpy_q_f32<<>> + cpy_q_f32<<>> (cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13); } @@ -276,9 +276,9 @@ static void ggml_cpy_f32_q4_0_cuda( const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) { GGML_ASSERT(ne % QK4_0 == 0); - const int64_t num_blocks = ne / QK4_0; + const int64_t num_blocks = (ne/QK4_0 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE; GGML_ASSERT(num_blocks <= INT_MAX); - cpy_f32_q<<>> + cpy_f32_q<<>> (cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13); } @@ -289,9 +289,9 @@ static void ggml_cpy_q4_0_f32_cuda( const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) { - const int64_t num_blocks = ne; + const int64_t num_blocks = (ne/QK4_0 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE; GGML_ASSERT(num_blocks <= INT_MAX); - cpy_q_f32, QK4_0><<>>( + cpy_q_f32, QK4_0><<>>( cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13); } @@ -302,9 +302,9 @@ static void ggml_cpy_f32_q4_1_cuda( const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) { GGML_ASSERT(ne % QK4_1 == 0); - const int64_t num_blocks = ne / QK4_1; + const int64_t num_blocks = (ne/QK4_1 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE; GGML_ASSERT(num_blocks <= INT_MAX); - cpy_f32_q<<>> + cpy_f32_q<<>> (cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13); } @@ -315,9 +315,9 @@ static void ggml_cpy_q4_1_f32_cuda( const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) { - const int64_t num_blocks = ne; + const int64_t num_blocks = (ne/QK4_1 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE; GGML_ASSERT(num_blocks <= INT_MAX); - cpy_q_f32, QK4_1><<>>( + cpy_q_f32, QK4_1><<>>( cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13); } @@ -328,9 +328,9 @@ static void ggml_cpy_f32_q5_0_cuda( const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) { GGML_ASSERT(ne % QK5_0 == 0); - const int64_t num_blocks = ne / QK5_0; + const int64_t num_blocks = (ne/QK5_0 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE; GGML_ASSERT(num_blocks <= INT_MAX); - cpy_f32_q<<>> + cpy_f32_q<<>> (cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13); } @@ -341,9 +341,9 @@ static void ggml_cpy_q5_0_f32_cuda( const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) { - const int64_t num_blocks = ne; + const int64_t num_blocks = (ne/QK5_0 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE; GGML_ASSERT(num_blocks <= INT_MAX); - cpy_q_f32, QK5_0><<>>( + cpy_q_f32, QK5_0><<>>( cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13); } @@ -354,9 +354,9 @@ static void ggml_cpy_f32_q5_1_cuda( const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) { GGML_ASSERT(ne % QK5_1 == 0); - const int64_t num_blocks = ne / QK5_1; + const int64_t num_blocks = (ne/QK5_1 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE; GGML_ASSERT(num_blocks <= INT_MAX); - cpy_f32_q<<>> + cpy_f32_q<<>> (cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13); } @@ -367,9 +367,9 @@ static void ggml_cpy_q5_1_f32_cuda( const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) { - const int64_t num_blocks = ne; + const int64_t num_blocks = (ne/QK5_1 + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE; GGML_ASSERT(num_blocks <= INT_MAX); - cpy_q_f32, QK5_1><<>>( + cpy_q_f32, QK5_1><<>>( cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13); } @@ -380,9 +380,9 @@ static void ggml_cpy_f32_iq4_nl_cuda( const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) { GGML_ASSERT(ne % QK4_NL == 0); - const int64_t num_blocks = ne / QK4_NL; + const int64_t num_blocks = (ne/QK4_NL + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE; GGML_ASSERT(num_blocks <= INT_MAX); - cpy_f32_q<<>> + cpy_f32_q<<>> (cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13); } diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 6d474bc11..6688debd4 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -8576,6 +8576,9 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_cpy(type_src, type_dst, {256, 2, 3, 4}, {-1,-1,-1,-1}, {1, 0, 2, 3})); // cpy not-contiguous } } + // quant block count not a multiple of the kernel block size + test_cases.emplace_back(new test_cpy(GGML_TYPE_F32, GGML_TYPE_Q4_0, {96, 1, 1, 1})); + test_cases.emplace_back(new test_cpy(GGML_TYPE_Q4_0, GGML_TYPE_F32, {96, 1, 1, 1})); test_cases.emplace_back(new test_cpy(GGML_TYPE_F32, GGML_TYPE_I32, {256, 2, 3, 4})); test_cases.emplace_back(new test_cpy(GGML_TYPE_F32, GGML_TYPE_I32, {256, 2, 3, 4}, {-1,-1,-1,-1}, {1, 0, 2, 3})); test_cases.emplace_back(new test_cpy(GGML_TYPE_I32, GGML_TYPE_F32, {256, 2, 3, 4})); From dd2c7c44710e860a428b46a92e2a9e39c428628b Mon Sep 17 00:00:00 2001 From: Xuan-Son Nguyen Date: Sat, 8 Aug 2026 16:35:53 +0200 Subject: [PATCH 02/16] server: add initial tool isolation support (via docker) (#26507) * server: add initial tool isolation support (via docker) * add docs * adapt get_info * py: fix type check * cont * separate tools_io_sandbox / tools_io_docker * rename sandbox --> isolate * x-tool-docker --> x-tool-runtime --------- Co-authored-by: Pascal --- common/arg.cpp | 10 + common/common.h | 1 + tools/cli/README.md | 1 - tools/completion/README.md | 1 - tools/server/README-dev.md | 1 + tools/server/README.md | 3 +- tools/server/server-tools.cpp | 567 +++++++++++++++--- tools/server/server-tools.h | 11 +- tools/server/server.cpp | 5 +- tools/server/tests/unit/test_tools_builtin.py | 91 +++ tools/server/tests/utils.py | 3 + 11 files changed, 613 insertions(+), 81 deletions(-) diff --git a/common/arg.cpp b/common/arg.cpp index da4087474..4cb853c7a 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -3308,6 +3308,16 @@ common_params_context common_params_parser_init(common_params & params, llama_ex params.server_tools = parse_csv_row(value); } ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_TOOLS")); + add_opt(common_arg( + {"--tools-runtime"}, "OPTION", + "experimental: run tools in a separate runtime environment (default: none, use host environment)\n" + "available options:\n" + " 'docker:': spin up a new Docker container and reuse it for all invocations, clean up on server exit\n" + " 'docker-container:': use an existing Docker container by ID, won't stop on server exit\n", + [](common_params & params, const std::string & value) { + params.server_tools_runtime = value; + } + ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_TOOLS_RUNTIME")); add_opt(common_arg( {"--mcp-servers-config"}, "PATH", "experimental: path to JSON file with MCP server definitions (Cursor-compatible format) - do not enable in untrusted environments (default: none)\n" diff --git a/common/common.h b/common/common.h index 2e15ec3f8..4811345f9 100644 --- a/common/common.h +++ b/common/common.h @@ -655,6 +655,7 @@ struct common_params { // enable built-in tools std::vector server_tools; + std::string server_tools_runtime; // MCP server configs (Cursor-compatible JSON) std::string mcp_servers_config; // path to JSON file with MCP server definitions diff --git a/tools/cli/README.md b/tools/cli/README.md index 4d86ce7c0..640d4fee8 100644 --- a/tools/cli/README.md +++ b/tools/cli/README.md @@ -54,7 +54,6 @@ | `-ctv, --cache-type-v TYPE` | KV cache data type for V
allowed values: f32, f16, bf16, q8_0, q4_0, q4_1, iq4_nl, q5_0, q5_1
(default: f16)
(env: LLAMA_ARG_CACHE_TYPE_V) | | `-dt, --defrag-thold N` | KV cache defragmentation threshold (DEPRECATED)
(env: LLAMA_ARG_DEFRAG_THOLD) | | `-np, --parallel N` | number of parallel sequences to decode (default: 1)
(env: LLAMA_ARG_N_PARALLEL) | -| `--rpc SERVERS` | comma-separated list of RPC servers (host:port)
(env: LLAMA_ARG_RPC) | | `--mlock` | DEPRECATED in favor of `--load-mode`: force system to keep model in RAM rather than swapping or compressing
(env: LLAMA_ARG_MLOCK) | | `--mmap, --no-mmap` | DEPRECATED in favor of `--load-mode`: whether to memory-map model. (if mmap disabled, slower load but may reduce pageouts if not using mlock)
(env: LLAMA_ARG_MMAP) | | `-dio, --direct-io, -ndio, --no-direct-io` | DEPRECATED in favor of `--load-mode`: use DirectIO if available
(env: LLAMA_ARG_DIO) | diff --git a/tools/completion/README.md b/tools/completion/README.md index 2abe7aaa2..e0923ea30 100644 --- a/tools/completion/README.md +++ b/tools/completion/README.md @@ -137,7 +137,6 @@ llama-completion.exe -m models\gemma-1.1-7b-it.Q4_K_M.gguf --ignore-eos -n -1 | `-ctv, --cache-type-v TYPE` | KV cache data type for V
allowed values: f32, f16, bf16, q8_0, q4_0, q4_1, iq4_nl, q5_0, q5_1
(default: f16)
(env: LLAMA_ARG_CACHE_TYPE_V) | | `-dt, --defrag-thold N` | KV cache defragmentation threshold (DEPRECATED)
(env: LLAMA_ARG_DEFRAG_THOLD) | | `-np, --parallel N` | number of parallel sequences to decode (default: 1)
(env: LLAMA_ARG_N_PARALLEL) | -| `--rpc SERVERS` | comma-separated list of RPC servers (host:port)
(env: LLAMA_ARG_RPC) | | `--mlock` | DEPRECATED in favor of `--load-mode`: force system to keep model in RAM rather than swapping or compressing
(env: LLAMA_ARG_MLOCK) | | `--mmap, --no-mmap` | DEPRECATED in favor of `--load-mode`: whether to memory-map model. (if mmap disabled, slower load but may reduce pageouts if not using mlock)
(env: LLAMA_ARG_MMAP) | | `-dio, --direct-io, -ndio, --no-direct-io` | DEPRECATED in favor of `--load-mode`: use DirectIO if available
(env: LLAMA_ARG_DIO) | diff --git a/tools/server/README-dev.md b/tools/server/README-dev.md index 45bcdcca7..31408f426 100644 --- a/tools/server/README-dev.md +++ b/tools/server/README-dev.md @@ -201,6 +201,7 @@ Invoke a tool call, request body is a JSON object with: Headers: - `x-tool-cwd`: optional; if set, use as the CWD for tool; this is not part of tool's params because it's meant to be set by the runtime, not the LLM itself +- `x-tool-runtime`: optional; if set, run the tool inside this isolate instead of on the host. Only `docker-container:` is supported for now, using an already-running container Returns JSON object. There are two response formats (MCP tools use the same two formats: their result content is concatenated into `plain_text_response`, and RPC or tool errors are surfaced as the `error` string): diff --git a/tools/server/README.md b/tools/server/README.md index 4d80f059d..64f0b0326 100644 --- a/tools/server/README.md +++ b/tools/server/README.md @@ -71,7 +71,6 @@ For the full list of features, please refer to [server's changelog](https://gith | `-ctk, --cache-type-k TYPE` | KV cache data type for K
allowed values: f32, f16, bf16, q8_0, q4_0, q4_1, iq4_nl, q5_0, q5_1
(default: f16)
(env: LLAMA_ARG_CACHE_TYPE_K) | | `-ctv, --cache-type-v TYPE` | KV cache data type for V
allowed values: f32, f16, bf16, q8_0, q4_0, q4_1, iq4_nl, q5_0, q5_1
(default: f16)
(env: LLAMA_ARG_CACHE_TYPE_V) | | `-dt, --defrag-thold N` | KV cache defragmentation threshold (DEPRECATED)
(env: LLAMA_ARG_DEFRAG_THOLD) | -| `--rpc SERVERS` | comma-separated list of RPC servers (host:port)
(env: LLAMA_ARG_RPC) | | `--mlock` | DEPRECATED in favor of `--load-mode`: force system to keep model in RAM rather than swapping or compressing
(env: LLAMA_ARG_MLOCK) | | `--mmap, --no-mmap` | DEPRECATED in favor of `--load-mode`: whether to memory-map model. (if mmap disabled, slower load but may reduce pageouts if not using mlock)
(env: LLAMA_ARG_MMAP) | | `-dio, --direct-io, -ndio, --no-direct-io` | DEPRECATED in favor of `--load-mode`: use DirectIO if available
(env: LLAMA_ARG_DIO) | @@ -198,6 +197,8 @@ For the full list of features, please refer to [server's changelog](https://gith | `--ui-config, --webui-config JSON` | JSON that provides default UI settings (overrides UI defaults)
(env: LLAMA_ARG_UI_CONFIG) | | `--ui-config-file, --webui-config-file PATH` | JSON file that provides default UI settings (overrides UI defaults)
(env: LLAMA_ARG_UI_CONFIG_FILE) | | `--ui-mcp-proxy, --webui-mcp-proxy, --no-ui-mcp-proxy, --no-webui-mcp-proxy` | experimental: whether to enable MCP CORS proxy - do not enable in untrusted environments (default: disabled)
(env: LLAMA_ARG_UI_MCP_PROXY) | +| `--tools TOOL1,TOOL2,...` | experimental: whether to enable built-in tools for AI agents - do not enable in untrusted environments (default: no tools)
specify "all" to enable all tools
available tools: read_file, file_glob_search, grep_search, exec_shell_command, write_file, edit_file, get_datetime
note: for security reasons, this will limit --cors-origins to localhost by default
(env: LLAMA_ARG_TOOLS) | +| `--tools-runtime OPTION` | experimental: run tools in a separate runtime environment (default: none, use host environment)
available options:
'docker:': spin up a new Docker container and reuse it for all invocations, clean up on server exit
'docker-container:': use an existing Docker container by ID, won't stop on server exit

(env: LLAMA_ARG_TOOLS_RUNTIME) | | `--tools TOOL1,TOOL2,...` | experimental: whether to enable built-in tools for AI agents - do not enable in untrusted environments (default: no tools)
specify "all" to enable all tools
available tools: read_file, file_glob_search, grep_search, exec_shell_command, write_file, edit_file, get_datetime, get_info
note: for security reasons, this will limit --cors-origins to localhost by default
(env: LLAMA_ARG_TOOLS) | | `--mcp-servers-config PATH` | experimental: path to JSON file with MCP server definitions (Cursor-compatible format) - do not enable in untrusted environments (default: none)
note: for security reasons, this will limit --cors-origins to localhost by default
(env: LLAMA_ARG_MCP_SERVERS_CONFIG) | | `--mcp-servers-json JSON` | experimental: inline JSON with MCP server definitions (Cursor-compatible format) - do not enable in untrusted environments (default: none)
note: for security reasons, this will limit --cors-origins to localhost by default
(env: LLAMA_ARG_MCP_SERVERS_JSON) | diff --git a/tools/server/server-tools.cpp b/tools/server/server-tools.cpp index 2fcb2a3c8..eacfbf0f7 100644 --- a/tools/server/server-tools.cpp +++ b/tools/server/server-tools.cpp @@ -10,12 +10,15 @@ #include #include #include +#include #include #include +#include #include #include #include #include +#include #if defined(_WIN32) # ifndef NOMINMAX @@ -127,6 +130,13 @@ static int entry_depth(const std::string & rel) { return 1 + (int) std::count(rel.begin(), rel.end(), '/'); } +// directories that a listing reports but never descends into: they can be enormous +// lowercase only, the local walker case-folds a name before the lookup +static const char * const SERVER_TOOL_JUNK_DIR_NAMES[] = { + ".git", ".svn", ".hg", "node_modules", "__pycache__", + ".venv", "venv", "dist", "build", "target", ".cache", ".idea", ".vscode", +}; + class tools_io { public: struct exec_result { @@ -165,6 +175,85 @@ public: const std::function & on_chunk = nullptr) const = 0; }; +// shared subprocess execution helper, used by both the local and the docker-backed tools_io implementations. +// combine_stderr=false when the raw stdout bytes must not be tainted by stderr, e.g. reading file contents. +static tools_io::exec_result run_subprocess( + const std::vector & args, + size_t max_output, + int timeout_secs, + const std::function & on_chunk, + bool combine_stderr, + const std::string & cwd = "") { + tools_io::exec_result res; + + common_subproc proc; + + int options = subprocess_option_no_window + | subprocess_option_inherit_environment + | subprocess_option_search_user_path; + if (combine_stderr) { + options |= subprocess_option_combined_stdout_stderr; + } + + if (!proc.create(args, options, {}, cwd.empty() ? nullptr : cwd.c_str())) { + res.output = "failed to spawn process"; + return res; + } + + std::atomic done{false}; + std::atomic timed_out{false}; + + std::thread timeout_thread([&]() { + auto deadline = std::chrono::steady_clock::now() + std::chrono::seconds(timeout_secs); + while (!done.load()) { + if (std::chrono::steady_clock::now() >= deadline) { + timed_out.store(true); + proc.terminate(); + return; + } + std::this_thread::sleep_for(std::chrono::milliseconds(100)); + } + }); + + FILE * f = proc.stdout_file(); + std::string output; + bool truncated = false; + if (f) { + char buf[4096]; + while (fgets(buf, sizeof(buf), f) != nullptr) { + if (!truncated) { + size_t len = strlen(buf); + if (output.size() + len <= max_output) { + output.append(buf, len); + if (on_chunk && !on_chunk(console_output_to_utf8(std::string(buf, len)))) { + proc.terminate(); + break; + } + } else { + size_t remaining = max_output - output.size(); + output.append(buf, remaining); + if (on_chunk && remaining > 0) on_chunk(console_output_to_utf8(std::string(buf, remaining))); + truncated = true; + } + } + } + } + + done.store(true); + if (timeout_thread.joinable()) { + timeout_thread.join(); + } + + res.exit_code = proc.join(); + + res.output = console_output_to_utf8(output); + res.timed_out = timed_out.load(); + if (truncated) { + res.output += "\n[output truncated]"; + } + return res; +} + class tools_io_basic : public tools_io { public: // cwd, if non-empty, is used to resolve relative paths and as the working directory for run() @@ -276,72 +365,7 @@ public: size_t max_output, int timeout_secs, const std::function & on_chunk = nullptr) const override { - exec_result res; - - common_subproc proc; - - int options = subprocess_option_no_window - | subprocess_option_combined_stdout_stderr - | subprocess_option_inherit_environment - | subprocess_option_search_user_path; - - if (!proc.create(args, options, {}, cwd.empty() ? nullptr : cwd.c_str())) { - res.output = "failed to spawn process"; - return res; - } - - std::atomic done{false}; - std::atomic timed_out{false}; - - std::thread timeout_thread([&]() { - auto deadline = std::chrono::steady_clock::now() + std::chrono::seconds(timeout_secs); - while (!done.load()) { - if (std::chrono::steady_clock::now() >= deadline) { - timed_out.store(true); - proc.terminate(); - return; - } - std::this_thread::sleep_for(std::chrono::milliseconds(100)); - } - }); - - FILE * f = proc.stdout_file(); - std::string output; - bool truncated = false; - if (f) { - char buf[4096]; - while (fgets(buf, sizeof(buf), f) != nullptr) { - if (!truncated) { - size_t len = strlen(buf); - if (output.size() + len <= max_output) { - output.append(buf, len); - if (on_chunk && !on_chunk(console_output_to_utf8(std::string(buf, len)))) { - proc.terminate(); - break; - } - } else { - size_t remaining = max_output - output.size(); - output.append(buf, remaining); - if (on_chunk && remaining > 0) on_chunk(console_output_to_utf8(std::string(buf, remaining))); - truncated = true; - } - } - } - } - - done.store(true); - if (timeout_thread.joinable()) { - timeout_thread.join(); - } - - res.exit_code = proc.join(); - - res.output = console_output_to_utf8(output); - res.timed_out = timed_out.load(); - if (truncated) { - res.output += "\n[output truncated]"; - } - return res; + return run_subprocess(args, max_output, timeout_secs, on_chunk, /*combine_stderr=*/true, cwd); } private: @@ -384,10 +408,8 @@ private: } static const std::unordered_set & junk_dir_names() { - static const std::unordered_set names = { - ".git", ".svn", ".hg", "node_modules", "__pycache__", - ".venv", "venv", "dist", "build", "target", ".cache", ".idea", ".vscode", - }; + static const std::unordered_set names( + std::begin(SERVER_TOOL_JUNK_DIR_NAMES), std::end(SERVER_TOOL_JUNK_DIR_NAMES)); return names; } @@ -450,9 +472,274 @@ private: } }; +// timeout for auxiliary isolate calls (stat/mkdir/ls/cp helpers); exec_shell_command uses its own +// caller-controlled timeout instead, enforced separately in run() +static constexpr int SERVER_TOOL_ISOLATE_EXEC_TIMEOUT = 15; // seconds +static constexpr size_t SERVER_TOOL_ISOLATE_READ_FILE_MAX_SIZE = 64 * 1024 * 1024; // 64 MB + +// runs every tools_io operation as a command inside an isolate: a container, a remote host, ... +// the isolate is created, mounted, and torn down externally by the caller +// it must provide a POSIX environment: sh, cat, wc, mkdir, dirname, find, timeout +class tools_io_isolate : public tools_io { +public: + // cwd, if non-empty, is used to resolve relative paths and as the working directory for run() + explicit tools_io_isolate(std::string cwd = "") : cwd(std::move(cwd)) {} + + // resolves `path` against `cwd` if `path` is relative and `cwd` is set; otherwise returns `path` unchanged. + // isolate paths are always POSIX-style ('/'), regardless of host OS. + std::string resolve(const std::string & path) const override { + if (cwd.empty() || (!path.empty() && path[0] == '/')) { + return path; + } + return cwd + "/" + path; + } + + bool is_directory(const std::string & path) const override { + return shell_test("-d", resolve(path)); + } + + bool is_regular_file(const std::string & path) const override { + return shell_test("-f", resolve(path)); + } + + bool file_size(const std::string & path, uintmax_t & out_size) const override { + auto res = exec({"sh", "-c", "wc -c < \"$1\"", "_", resolve(path)}, 64, true); + if (res.exit_code != 0 || res.timed_out) return false; + try { + size_t pos; + out_size = (uintmax_t) std::stoull(res.output, &pos); + } catch (...) { + return false; + } + return true; + } + + bool read_file(const std::string & path, std::string & out) const override { + // combine_stderr=false: stderr must not be spliced into raw file bytes + auto res = exec({"cat", "--", resolve(path)}, SERVER_TOOL_ISOLATE_READ_FILE_MAX_SIZE, false); + if (res.exit_code != 0 || res.timed_out) return false; + out = res.output; + return true; + } + + bool write_file(const std::string & path, const std::string & content) const override { + std::string abs_path = resolve(path); + + std::error_code ec; + fs::path tmp_dir = fs::temp_directory_path(ec); + if (ec) return false; + + static std::atomic tmp_counter{0}; + fs::path tmp = tmp_dir / string_format( + "llama-tools-io-isolate-%zu-%llu.tmp", + std::hash{}(std::this_thread::get_id()), + (unsigned long long) tmp_counter.fetch_add(1)); + + { + std::ofstream f(tmp, std::ios::binary); + if (!f) return false; + f << content; + if (!f) return false; + } + + bool ok = shell_run({"sh", "-c", "mkdir -p \"$(dirname \"$1\")\"", "_", abs_path}); + if (ok) { + ok = upload(tmp.string(), abs_path); + } + + std::error_code rm_ec; + fs::remove(tmp, rm_ec); + return ok; + } + + list_result list_entries(const std::string & base, int max_depth, list_kind kind) const override { + list_result out; + + const std::string abs_base = resolve(base); + if (!is_directory(base)) { + out.err = "path does not exist or is not a directory"; + return out; + } + + // git ls-files cannot list directories; use the walker when they are requested + if (kind == list_kind::files) { + auto res = exec( + {"sh", "-c", "cd \"$1\" && git ls-files --cached --others --exclude-standard", "_", abs_base}, + SERVER_TOOL_GIT_LS_FILES_MAX_OUTPUT, true); + + if (res.exit_code == 0 && !res.timed_out) { + for (const auto & rel : split_lines(res.output, /*strip_dot_slash=*/false)) { + if (max_depth > 0 && entry_depth(rel) > max_depth) continue; + out.entries.push_back({rel, false}); + } + return out; + } + } + + if (kind == list_kind::dirs || kind == list_kind::all) { + for (auto & rel : find_entries(abs_base, max_depth, /*dirs=*/true, out.truncated)) { + out.entries.push_back({std::move(rel), true}); + } + } + if (kind == list_kind::files || kind == list_kind::all) { + for (auto & rel : find_entries(abs_base, max_depth, /*dirs=*/false, out.truncated)) { + out.entries.push_back({std::move(rel), false}); + } + } + + return out; + } + + // wraps the command with an in-isolate `timeout`, since killing the host-side client + // does not kill the process tree running inside the isolate + exec_result run( + const std::vector & args, + size_t max_output, + int timeout_secs, + const std::function & on_chunk = nullptr) const override { + std::vector inner = {"timeout", std::to_string(timeout_secs) + "s"}; + inner.insert(inner.end(), args.begin(), args.end()); + // small buffer over timeout_secs so the in-isolate `timeout` has a chance to exit cleanly + // before the host-side supervisory timeout forcibly kills the client + return run_subprocess( + build_argv(with_cwd(inner), /*needs_stdin=*/true), + max_output, timeout_secs + 5, on_chunk, true); + } + +protected: + // wrap `inner` (a complete POSIX argv) into the host-side argv that runs it in the isolate + // a transport that re-parses its args in a remote shell (ssh) must join `inner` with shell_quote_join() + virtual std::vector build_argv(const std::vector & inner, bool needs_stdin) const = 0; + + // copy a host file into the isolate, `isolate_path` is absolute and its parent already exists + virtual bool upload(const std::string & host_path, const std::string & isolate_path) const = 0; + + // quote `argv` into a single string that a POSIX shell re-parses into exactly `argv` + static std::string shell_quote_join(const std::vector & argv) { + std::string out; + for (const auto & arg : argv) { + if (!out.empty()) out += ' '; + out += '\''; + for (const char c : arg) { + // a single quote cannot be escaped inside single quotes: close, escape, reopen + if (c == '\'') out += "'\\''"; + else out += c; + } + out += '\''; + } + return out; + } + +private: + std::string cwd; + + // set the working directory in the command itself, docker's `-w` has no equivalent on every transport + // auxiliary calls do not need this, they use the absolute paths from resolve() + std::vector with_cwd(const std::vector & inner) const { + if (cwd.empty()) { + return inner; + } + // 127 is what a shell reports for a command it could not run + std::vector out = {"sh", "-c", "cd \"$1\" || exit 127; shift; exec \"$@\"", "_", cwd}; + out.insert(out.end(), inner.begin(), inner.end()); + return out; + } + + exec_result exec(const std::vector & inner, size_t max_output, bool combine_stderr) const { + return run_subprocess( + build_argv(inner, /*needs_stdin=*/false), + max_output, SERVER_TOOL_ISOLATE_EXEC_TIMEOUT, nullptr, combine_stderr); + } + + bool shell_run(const std::vector & inner) const { + auto res = exec(inner, 4096, true); + return res.exit_code == 0 && !res.timed_out; + } + + bool shell_test(const char * flag, const std::string & path) const { + return shell_run({"sh", "-c", std::string("[ ") + flag + " \"$1\" ]", "_", path}); + } + + static std::vector split_lines(const std::string & text, bool strip_dot_slash) { + std::vector result; + std::istringstream iss(text); + std::string line; + while (std::getline(iss, line)) { + if (!line.empty() && line.back() == '\r') line.pop_back(); + if (line.empty()) continue; + if (strip_dot_slash && line.rfind("./", 0) == 0) line = line.substr(2); + std::replace(line.begin(), line.end(), '\\', '/'); + result.push_back(line); + } + return result; + } + + // one `find` pass in the isolate. junk directories stay selectable but are never descended into, + // and -mindepth/-maxdepth keep a busybox image working as well as a GNU one + std::vector find_entries(const std::string & abs_base, int max_depth, bool dirs, bool & truncated) const { + std::string prune_expr; + for (const char * n : SERVER_TOOL_JUNK_DIR_NAMES) { + if (!prune_expr.empty()) prune_expr += " -o "; + prune_expr += std::string("-name ") + n; + } + + std::string cmd = "cd \"$1\" && find . -mindepth 1"; + if (max_depth > 0) { + cmd += " -maxdepth " + std::to_string(max_depth); + } + cmd += " \\( " + prune_expr + " \\) -prune"; + cmd += dirs ? " -print -o -type d -print" : " -o -type f -print"; + + auto res = exec({"sh", "-c", cmd, "_", abs_base}, SERVER_TOOL_GIT_LS_FILES_MAX_OUTPUT, true); + truncated = truncated || res.timed_out; + return split_lines(res.output, /*strip_dot_slash=*/true); + } +}; + +// an already-running docker container, driven through `docker exec` and `docker cp` +class tools_io_docker : public tools_io_isolate { +public: + tools_io_docker(std::string container_id, std::string cwd = "") + : tools_io_isolate(std::move(cwd)), container_id(std::move(container_id)) {} + +protected: + std::vector build_argv(const std::vector & inner, bool needs_stdin) const override { + std::vector argv = {"docker", "exec"}; + if (needs_stdin) { + argv.push_back("-i"); + } + argv.push_back(container_id); + argv.insert(argv.end(), inner.begin(), inner.end()); + return argv; + } + + bool upload(const std::string & host_path, const std::string & isolate_path) const override { + auto res = run_subprocess( + {"docker", "cp", host_path, container_id + ":" + isolate_path}, + 4096, SERVER_TOOL_ISOLATE_EXEC_TIMEOUT, nullptr, true); + return res.exit_code == 0 && !res.timed_out; + } + +private: + std::string container_id; +}; + +// runtime spec used by --tools-runtime and the x-tool-runtime header +// this is the only scheme for now, ssh: and podman: can be added next to it +static const std::string SERVER_TOOL_RUNTIME_DOCKER_CONTAINER = "docker-container:"; + +// an empty runtime runs the tools on the host static std::unique_ptr make_tools_io(const json & params) { - std::string cwd = json_value(params, "cwd", std::string()); - return std::make_unique(cwd); + std::string cwd = json_value(params, "cwd", std::string()); + std::string runtime = json_value(params, "runtime", std::string()); + if (runtime.empty()) { + return std::make_unique(cwd); + } + if (runtime.rfind(SERVER_TOOL_RUNTIME_DOCKER_CONTAINER, 0) == 0) { + return std::make_unique(runtime.substr(SERVER_TOOL_RUNTIME_DOCKER_CONTAINER.size()), cwd); + } + // do not fall back to the host, the caller asked for an isolate + throw std::runtime_error("unknown tool runtime: " + runtime); } // no '/' in pattern -> match basename at any depth; else match full relative path @@ -861,8 +1148,11 @@ struct server_tool_exec_shell_command : server_tool { timeout = std::min(timeout, SERVER_TOOL_EXEC_SHELL_COMMAND_MAX_TIMEOUT); max_output = std::min(max_output, SERVER_TOOL_EXEC_SHELL_COMMAND_MAX_OUTPUT_SIZE); + // an isolate is always POSIX regardless of host OS, so it always gets `sh -c` #ifdef _WIN32 - std::vector args = {"cmd", "/c", command}; + std::vector args = !json_value(params, "runtime", std::string()).empty() + ? std::vector{"sh", "-c", command} + : std::vector{"cmd", "/c", command}; #else std::vector args = {"sh", "-c", command}; #endif @@ -1355,11 +1645,16 @@ struct server_tool_get_info : server_tool { json invoke(json params, server_tool::stream *) const override { auto io = make_tools_io(params); + // inside an isolate, we always use the linux command #ifdef _WIN32 - auto res = io->run({"cmd", "/c", "ver"}, SERVER_TOOL_GET_INFO_MAX_OUTPUT, SERVER_TOOL_GET_INFO_TIMEOUT); + std::vector args = !json_value(params, "runtime", std::string()).empty() + ? std::vector{"uname", "-a"} + : std::vector{"cmd", "/c", "ver"}; #else - auto res = io->run({"uname", "-a"}, SERVER_TOOL_GET_INFO_MAX_OUTPUT, SERVER_TOOL_GET_INFO_TIMEOUT); + std::vector args = {"uname", "-a"}; #endif + + auto res = io->run(args, SERVER_TOOL_GET_INFO_MAX_OUTPUT, SERVER_TOOL_GET_INFO_TIMEOUT); // "ver" prints a blank line before the version, so the output is stripped on both ends; // a failed spawn or a timeout leaves a diagnostic in res.output, which is not an OS name std::string os_info = res.exit_code == 0 && !res.timed_out ? string_strip(res.output) : "unknown"; @@ -1461,6 +1756,103 @@ struct server_mcp_tool : server_tool { } }; +// owns the docker container used as the sandboxed runtime for tool invocations, as configured by +// --tools-runtime. "spawned" mode starts and stops the container itself; "existing" mode just reuses +// a container id the user already has running and never stops it. +struct server_tools_docker_runtime { + server_tools_docker_runtime(const server_tools_docker_runtime &) = delete; + + explicit server_tools_docker_runtime(const std::string & spec) { + static const std::string docker_prefix = "docker:"; + if (spec.rfind(docker_prefix, 0) == 0) { + spawned = true; + image = spec.substr(docker_prefix.size()); + if (image.empty()) { + throw std::runtime_error("--tools-runtime docker: requires an image name"); + } + spawn(); + } else if (spec.rfind(SERVER_TOOL_RUNTIME_DOCKER_CONTAINER, 0) == 0) { + spawned = false; + container_id = spec.substr(SERVER_TOOL_RUNTIME_DOCKER_CONTAINER.size()); + if (container_id.empty()) { + throw std::runtime_error("--tools-runtime docker-container: requires a container id"); + } + } else { + throw std::runtime_error("unknown --tools-runtime option: " + spec); + } + } + + ~server_tools_docker_runtime() { + if (spawned && !container_id.empty()) { + // closing stdin signals the container's shell (its pid 1) to exit; --rm then removes it + proc.close_stdin(); + proc.join(); + } + } + + // container id to use for the next tool call; respawns a spawned container that died on its own, + // or throws if an externally-managed one is no longer reachable + std::string get_container_id() { + std::lock_guard lock(mutex); + if (!spawned) { + if (!is_running(container_id)) { + throw std::runtime_error(string_format( + "docker container \"%s\" is no longer running, restart it to keep using tools", + container_id.c_str())); + } + return container_id; + } + + if (!proc.alive()) { + SRV_WRN("docker tools runtime container \"%s\" died, respawning\n", container_id.c_str()); + spawn(); + } + return container_id; + } + +private: + bool spawned = false; + std::string image; // spawned mode only + std::string container_id; + common_subproc proc; // spawned mode only: `docker run` client that keeps the container alive + std::mutex mutex; + + // spawns "docker run --rm -i sh" and keeps its stdin open; the shell blocks reading stdin, + // so the container stays alive until we close it (see destructor) or it is killed from the outside + void spawn() { + std::error_code ec; + fs::path cidfile = fs::temp_directory_path(ec) / string_format( + "llama-tools-runtime-cid-%zu.tmp", std::hash{}(std::this_thread::get_id())); + fs::remove(cidfile, ec); + + std::vector args = {"docker", "run", "--rm", "-i", "--cidfile", cidfile.string(), image, "sh"}; + int options = subprocess_option_no_window + | subprocess_option_inherit_environment + | subprocess_option_search_user_path; + if (!proc.create(args, options)) { + throw std::runtime_error("failed to spawn docker container for tools runtime (image: " + image + ")"); + } + + std::string cid; + for (int i = 0; i < 100 && cid.empty(); i++) { + std::ifstream f(cidfile); + if (f) std::getline(f, cid); + if (cid.empty()) std::this_thread::sleep_for(std::chrono::milliseconds(100)); + } + fs::remove(cidfile, ec); + if (cid.empty()) { + proc.terminate(); + throw std::runtime_error("timed out waiting for docker container to start (image: " + image + ")"); + } + container_id = cid; + } + + static bool is_running(const std::string & id) { + auto res = run_subprocess({"docker", "inspect", "-f", "{{.State.Running}}", id}, 16, 5, nullptr, true); + return res.exit_code == 0 && !res.timed_out && res.output.rfind("true", 0) == 0; + } +}; + static server_tool & find_tool(std::vector> & tools, const std::string & name, bool require_stream) { for (auto & t : tools) { if (t->name == name) { @@ -1506,8 +1898,16 @@ static std::string get_header(const std::map & headers return default_value; } +server_tools::server_tools() = default; +server_tools::~server_tools() = default; + void server_tools::setup(const std::vector & enabled_tools, - server_mcp & mcp_mgr) { + server_mcp & mcp_mgr, + const std::string & tools_runtime) { + if (!tools_runtime.empty()) { + docker_runtime = std::make_unique(tools_runtime); + } + if (!enabled_tools.empty()) { if (!common_subproc::is_supported()) { throw std::runtime_error("subprocess is not enabled on this build"); @@ -1590,11 +1990,26 @@ void server_tools::setup(const std::vector & enabled_tools, bool stream = body.value("stream", false); // accept x-tool-cwd header to override of the process + if (params.contains("cwd")) { + params.erase("cwd"); + } auto cwd = get_header(req.headers, "x-tool-cwd"); if (!cwd.empty()) { params["cwd"] = cwd; } + // accept x-tool-runtime header to route tool I/O through an isolate, e.g. "docker-container:"; + // falls back to the --tools-runtime isolate, if configured + if (params.contains("runtime")) { + params.erase("runtime"); + } + auto runtime = get_header(req.headers, "x-tool-runtime"); + if (!runtime.empty()) { + params["runtime"] = runtime; + } else if (docker_runtime) { + params["runtime"] = SERVER_TOOL_RUNTIME_DOCKER_CONTAINER + docker_runtime->get_container_id(); + } + server_tool & tool = find_tool(tools, tool_name, stream); if (stream) { diff --git a/tools/server/server-tools.h b/tools/server/server-tools.h index 601399ee9..7f70e6767 100644 --- a/tools/server/server-tools.h +++ b/tools/server/server-tools.h @@ -30,6 +30,8 @@ struct server_tool { json to_json() const; }; +struct server_tools_docker_runtime; // impl detail, defined in server-tools.cpp + struct server_tools { std::vector> tools; @@ -37,9 +39,16 @@ struct server_tools { server_response queue_res; std::atomic res_id{0}; + // set when --tools-runtime is configured; owns the docker container used to run tools, if any + std::unique_ptr docker_runtime; + void setup(const std::vector & enabled_tools, - server_mcp & mcp_mgr); + server_mcp & mcp_mgr, + const std::string & tools_runtime); server_http_context::handler_t handle_get; server_http_context::handler_t handle_post; + + server_tools(); + ~server_tools(); }; diff --git a/tools/server/server.cpp b/tools/server/server.cpp index aafb1f307..1b2e6edb4 100644 --- a/tools/server/server.cpp +++ b/tools/server/server.cpp @@ -338,7 +338,7 @@ int llama_server(common_params & params, int argc, char ** argv) { if (!params.server_tools.empty() || !mcp_mgr.empty()) { try { - tools.setup(params.server_tools, mcp_mgr); + tools.setup(params.server_tools, mcp_mgr, params.server_tools_runtime); } catch (const std::exception & e) { SRV_ERR("tools setup failed: %s\n", e.what()); return 1; @@ -348,6 +348,9 @@ int llama_server(common_params & params, int argc, char ** argv) { if (!params.server_tools.empty()) { warn_names.push_back("built-in tools (experimental)"); } + if (!params.server_tools_runtime.empty()) { + warn_names.push_back("tools runtime (experimental)"); + } if (!mcp_mgr.empty()) { warn_names.push_back("MCP servers (experimental)"); } diff --git a/tools/server/tests/unit/test_tools_builtin.py b/tools/server/tests/unit/test_tools_builtin.py index 11c82e690..c651e8e72 100755 --- a/tools/server/tests/unit/test_tools_builtin.py +++ b/tools/server/tests/unit/test_tools_builtin.py @@ -1,4 +1,6 @@ import os +import shutil +import subprocess import pytest from utils import * @@ -146,6 +148,95 @@ def test_tools_builtin_cwd_header(): os.remove(marker_path) +def _docker_unavailable_reason() -> str | None: + """None if docker can be used to run a container, otherwise the reason it can't.""" + docker_bin = shutil.which("docker") + if docker_bin is None: + return "docker is not installed" + try: + subprocess.run([docker_bin, "info"], capture_output=True, timeout=5, check=True) + except Exception as e: + return f"docker daemon is not usable: {e}" + return None + + +@pytest.fixture +def docker_container(): + reason = _docker_unavailable_reason() + if reason is not None: + pytest.skip(reason) # ty: ignore[too-many-positional-arguments, invalid-argument-type] + + proc = subprocess.run( + ["docker", "run", "-d", "--rm", "busybox", "sleep", "300"], + capture_output=True, text=True, + ) + if proc.returncode != 0: + pytest.skip(f"failed to start docker container: {proc.stderr.strip()}") # ty: ignore[too-many-positional-arguments, invalid-argument-type] + + container_id = proc.stdout.strip() + try: + yield container_id + finally: + subprocess.run(["docker", "rm", "-f", container_id], capture_output=True) + + +def test_tools_builtin_runtime_header(docker_container: str): + global server + server.start() + + headers = {"x-tool-runtime": f"docker-container:{docker_container}", "x-tool-cwd": "/tmp"} + + write_res = call_tool("write_file", {"path": "test.log", "content": "hello docker\n"}, headers=headers) + assert write_res["result"] == "file written successfully" + + read_res = call_tool("read_file", {"path": "test.log"}, headers=headers) + assert read_res["plain_text_response"] == "hello docker\n" + + exec_res = call_tool("exec_shell_command", {"command": "cat test.log"}, headers=headers) + assert "hello docker" in exec_res["plain_text_response"] + + +def test_tools_builtin_runtime_header_unknown_scheme(): + global server + server.start() + + # an unknown runtime must fail, never silently fall back to running on the host + res = server.make_request("POST", "/tools", + data={"tool": "exec_shell_command", "params": {"command": "echo hi"}}, + headers={"x-tool-runtime": "ssh:example.com"}) + assert res.status_code == 500, res.body + assert "unknown tool runtime" in str(res.body) + + +def test_tools_builtin_docker_runtime_cleans_up_spawned_container(): + reason = _docker_unavailable_reason() + if reason is not None: + pytest.skip(reason) # ty: ignore[too-many-positional-arguments, invalid-argument-type] + + global server + server.server_tools_runtime = "docker:busybox" + server.start() + + # exec_shell_command runs inside the container spawned for --tools-runtime; docker sets + # the container's hostname to its own short id, so this also tells us which one to check + res = call_tool("exec_shell_command", {"command": "hostname"}) + container_id = res["plain_text_response"].splitlines()[0].strip() + assert len(container_id) >= 8, res + + running = subprocess.run( + ["docker", "inspect", "-f", "{{.State.Running}}", container_id], + capture_output=True, text=True, + ) + assert running.returncode == 0 and running.stdout.strip() == "true", running.stderr + + server.stop() + + # a clean server shutdown must stop and remove the container it spawned (it runs with --rm), + # not leave it behind as an abandoned child + leftover = subprocess.run(["docker", "inspect", container_id], capture_output=True, text=True) + assert leftover.returncode != 0, f"container {container_id} was not cleaned up after server exit" + + def test_tools_builtin_edit_file_rejects_overlapping_edits(): global server server.start() diff --git a/tools/server/tests/utils.py b/tools/server/tests/utils.py index 3416f0bfd..fffe07a67 100644 --- a/tools/server/tests/utils.py +++ b/tools/server/tests/utils.py @@ -115,6 +115,7 @@ class ServerProcess: backend_sampling: bool = False gcp_compat: bool = False server_tools: str | None = None + server_tools_runtime: str | None = None mcp_servers_config: str | None = None mcp_servers_json: str | None = None cors_origins: str | None = None @@ -270,6 +271,8 @@ class ServerProcess: server_args.append("--ui-mcp-proxy") if self.server_tools: server_args.extend(["--tools", self.server_tools]) + if self.server_tools_runtime: + server_args.extend(["--tools-runtime", self.server_tools_runtime]) if self.mcp_servers_config: server_args.extend(["--mcp-servers-config", self.mcp_servers_config]) if self.mcp_servers_json: From 18f7ad7fc912444acc0f51995a4b8e45fd9a0cd4 Mon Sep 17 00:00:00 2001 From: Pascal Date: Sat, 8 Aug 2026 16:36:21 +0200 Subject: [PATCH 03/16] server, ui: only offer a working directory when a tool reads it (#26762) The working directory chip showed up as soon as the server exposed any builtin tool, so a server started with just get_datetime, or a user who turned every filesystem tool off in the settings, still got a control that nothing would read. Tools now declare whether they resolve their paths and run against the working directory, next to the write permission they already publish in the /tools listing. The WebUI shows the chip and enables the /cwd command only when at least one such tool is both served and left enabled. --- tools/server/server-tools.cpp | 8 +++++++ tools/server/server-tools.h | 1 + .../app/chat/ChatForm/ChatForm.svelte | 4 ++-- tools/ui/src/lib/constants/chat-commands.ts | 4 ++-- .../lib/hooks/use-chat-form-pickers.svelte.ts | 4 ++-- tools/ui/src/lib/stores/tools.svelte.ts | 21 +++++++++++++++++++ tools/ui/src/lib/types/mcp.d.ts | 1 + .../components/ChatFormPickersHarness.svelte | 2 +- 8 files changed, 38 insertions(+), 7 deletions(-) diff --git a/tools/server/server-tools.cpp b/tools/server/server-tools.cpp index eacfbf0f7..d5e696434 100644 --- a/tools/server/server-tools.cpp +++ b/tools/server/server-tools.cpp @@ -74,6 +74,7 @@ json server_tool::to_json() const { {"permissions", json{ {"write", permission_write} }}, + {"uses_cwd", uses_cwd}, {"definition", get_definition()}, }; } @@ -763,6 +764,7 @@ struct server_tool_read_file : server_tool { server_tool_read_file() { name = "read_file"; display_name = "Read file"; + uses_cwd = true; permission_write = false; } @@ -851,6 +853,7 @@ struct server_tool_file_glob_search : server_tool { server_tool_file_glob_search() { name = "file_glob_search"; display_name = "File search"; + uses_cwd = true; permission_write = false; } @@ -965,6 +968,7 @@ struct server_tool_grep_search : server_tool { server_tool_grep_search() { name = "grep_search"; display_name = "Grep search"; + uses_cwd = true; permission_write = false; } @@ -1117,6 +1121,7 @@ struct server_tool_exec_shell_command : server_tool { server_tool_exec_shell_command() { name = "exec_shell_command"; display_name = "Execute shell command"; + uses_cwd = true; permission_write = true; support_stream = true; } @@ -1195,6 +1200,7 @@ struct server_tool_write_file : server_tool { server_tool_write_file() { name = "write_file"; display_name = "Write file"; + uses_cwd = true; permission_write = true; } @@ -1237,6 +1243,7 @@ struct server_tool_edit_file : server_tool { server_tool_edit_file() { name = "edit_file"; display_name = "Edit file"; + uses_cwd = true; permission_write = true; } @@ -1625,6 +1632,7 @@ struct server_tool_get_info : server_tool { server_tool_get_info() { name = "get_info"; display_name = "Get Runtime Info"; + uses_cwd = true; permission_write = false; } diff --git a/tools/server/server-tools.h b/tools/server/server-tools.h index 7f70e6767..ede303181 100644 --- a/tools/server/server-tools.h +++ b/tools/server/server-tools.h @@ -14,6 +14,7 @@ struct server_tool { std::string display_name; bool permission_write = false; bool support_stream = false; // if true, output can be streamed + bool uses_cwd = false; // if true, the tool resolves paths and runs against the working directory virtual ~server_tool() = default; virtual json get_definition() const = 0; diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatForm.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatForm.svelte index 1df125708..2d70c302c 100644 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatForm.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatForm.svelte @@ -156,7 +156,7 @@ focusInput: refocusInput, getShowModelSelector: () => showModelSelector, hasPrompts: () => mcpStore.hasPromptsCapability(conversationsStore.getAllMcpServerOverrides()), - hasBuiltinTools: () => toolsStore.builtinTools.length > 0, + hasCwdTools: () => toolsStore.hasEnabledCwdTools, getCwd: () => cwd, getServerHome: () => toolsStore.serverHome ?? null, openModelSelector: () => chatFormActionsRef?.openModelSelector(), @@ -651,7 +651,7 @@ - {#if toolsStore.builtinTools.length > 0} + {#if toolsStore.hasEnabledCwdTools} boolean; /** Gates `/cwd`. */ - hasBuiltinTools: () => boolean; + hasCwdTools: () => boolean; } /** @@ -32,7 +32,7 @@ export function getChatCommands(options: ChatCommandsOptions): ChatFormCommand[] description: SET_WORKING_DIRECTORY_LABEL, keywords: ['current working directory'], action: ChatFormCommandAction.CWD, - disabled: !options.hasBuiltinTools() + disabled: !options.hasCwdTools() }, { name: 'model', diff --git a/tools/ui/src/lib/hooks/use-chat-form-pickers.svelte.ts b/tools/ui/src/lib/hooks/use-chat-form-pickers.svelte.ts index f3bc5f732..4860f16fb 100644 --- a/tools/ui/src/lib/hooks/use-chat-form-pickers.svelte.ts +++ b/tools/ui/src/lib/hooks/use-chat-form-pickers.svelte.ts @@ -24,7 +24,7 @@ export interface UseChatFormPickersOptions { /** Gates `/prompt`. */ hasPrompts: () => boolean; /** Gates `/cwd`. */ - hasBuiltinTools: () => boolean; + hasCwdTools: () => boolean; getCwd: () => string | null; /** Mention search fallback scope. */ getServerHome: () => string | null; @@ -63,7 +63,7 @@ export function useChatFormPickers(opts: UseChatFormPickersOptions) { getChatCommands({ showModelSelector: opts.getShowModelSelector(), hasPrompts: opts.hasPrompts, - hasBuiltinTools: opts.hasBuiltinTools + hasCwdTools: opts.hasCwdTools }) ); diff --git a/tools/ui/src/lib/stores/tools.svelte.ts b/tools/ui/src/lib/stores/tools.svelte.ts index 5136101e7..4114ef756 100644 --- a/tools/ui/src/lib/stores/tools.svelte.ts +++ b/tools/ui/src/lib/stores/tools.svelte.ts @@ -27,6 +27,9 @@ class ToolsStore { private _loading = $state(false); private _error = $state(null); private _disabledTools = $state(new SvelteSet()); + // builtin tools that resolve their paths against the working directory, + // as declared by the server in its `/tools` listing + private _cwdAwareTools = $state(new SvelteSet()); private _toolsEndpointUnreachable = $state(false); private _serverHome = $state(undefined); @@ -476,6 +479,21 @@ class ToolsStore { return this.getEnabledToolsForLLM().length > 0; } + /** + * Check if a working directory is worth setting: at least one builtin tool + * that reads it is both served and left enabled by the user. + */ + get hasEnabledCwdTools(): boolean { + return this._builtinTools.some((def) => { + const name = def.function.name; + + return ( + this._cwdAwareTools.has(name) && + !this._disabledTools.has(this.toolKey(ToolSource.BUILTIN, name)) + ); + }); + } + async fetchBuiltinTools(): Promise { if (this._loading) return; @@ -486,6 +504,9 @@ class ToolsStore { try { const toolInfos = await ToolsService.list(); this._builtinTools = toolInfos.map((info) => info.definition); + this._cwdAwareTools = new SvelteSet( + toolInfos.filter((info) => info.uses_cwd).map((info) => info.tool) + ); } catch (err) { const errorMessage = err instanceof Error ? err.message : String(err); this._error = errorMessage; diff --git a/tools/ui/src/lib/types/mcp.d.ts b/tools/ui/src/lib/types/mcp.d.ts index b567c20c9..2b5937913 100644 --- a/tools/ui/src/lib/types/mcp.d.ts +++ b/tools/ui/src/lib/types/mcp.d.ts @@ -292,6 +292,7 @@ export interface ServerBuiltinToolInfo { permissions: { write: boolean; }; + uses_cwd: boolean; definition: OpenAIToolDefinition; } diff --git a/tools/ui/tests/client/components/ChatFormPickersHarness.svelte b/tools/ui/tests/client/components/ChatFormPickersHarness.svelte index 76b2e9fe3..8e5e56c20 100644 --- a/tools/ui/tests/client/components/ChatFormPickersHarness.svelte +++ b/tools/ui/tests/client/components/ChatFormPickersHarness.svelte @@ -21,7 +21,7 @@ focusInput: () => {}, getShowModelSelector: () => true, hasPrompts: () => true, - hasBuiltinTools: () => true, + hasCwdTools: () => true, getCwd: () => null, getServerHome: () => null, openModelSelector: () => { From 687e7789271ec1276e3470f158428e11a4f80b6f Mon Sep 17 00:00:00 2001 From: Rafail Giavrimis <47496212+grafail@users.noreply.github.com> Date: Sat, 8 Aug 2026 17:32:37 +0100 Subject: [PATCH 04/16] CUDA: fuse rms_norm + mul + rope (+ view + set_rows) (#26767) * CUDA: fuse rms_norm + mul + rope (+ view + set_rows) * tests: add broadcast weight case to rms_norm_mul_rope * CUDA: check memory ranges before rms_norm rope fusion * CUDA: check memory ranges in rope set_rows fusion --- ggml/src/ggml-cuda/ggml-cuda.cu | 89 +++++++++++- ggml/src/ggml-cuda/rope.cu | 235 ++++++++++++++++++++++++++++++++ ggml/src/ggml-cuda/rope.cuh | 2 + tests/test-backend-ops.cpp | 33 +++-- 4 files changed, 344 insertions(+), 15 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 5446b3131..dec619324 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -2651,6 +2651,52 @@ static bool ggml_cuda_should_fuse_rope_set_rows(const ggml_tensor * rope, return true; } +static bool ggml_cuda_should_fuse_rms_norm_mul_rope(const ggml_tensor * rms_norm, + const ggml_tensor * mul, + const ggml_tensor * rope) { + if (rms_norm->op != GGML_OP_RMS_NORM || mul->op != GGML_OP_MUL || rope->op != GGML_OP_ROPE) { + return false; + } + + if (rms_norm->src[0]->type != GGML_TYPE_F32 || rms_norm->type != GGML_TYPE_F32 || + mul->src[0]->type != GGML_TYPE_F32 || mul->src[1]->type != GGML_TYPE_F32 || + mul->type != GGML_TYPE_F32 || rope->type != GGML_TYPE_F32) { + return false; + } + + if (rope->src[0] != mul) { + return false; + } + + //if rms norm is the B operand, then we don't handle broadcast + if (rms_norm == mul->src[1] && !ggml_are_same_shape(mul->src[0], rms_norm)) { + return false; + } + + if (!ggml_are_same_shape(rms_norm, mul)) { + return false; + } + + //rms_norm kernel assumes contiguous rows + if (!ggml_is_contiguous_rows(rms_norm->src[0]) || + !ggml_is_contiguous_rows(mul->src[0]) || !ggml_is_contiguous_rows(mul->src[1])) { + return false; + } + + // the fused kernel handles the norm/neox rope modes only + const int mode = ((const int32_t *) rope->op_params)[2]; + if (mode != GGML_ROPE_TYPE_NORMAL && mode != GGML_ROPE_TYPE_NEOX) { + return false; + } + + const int n_dims = ((const int32_t *) rope->op_params)[1]; + if (n_dims % 2 != 0 || rope->src[0]->ne[0] % 2 != 0) { + return false; + } + + return true; +} + // match gated_delta_net + the strided cpy that scatters its state snapshots into the cache // (slot i -> rollback group i, slot 0 newest), so the kernel can write them and skip the cpy. static int ggml_cuda_try_gdn_cache_fusion( @@ -2980,6 +3026,36 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, } } + std::initializer_list rms_norm_mul_rope_ops = { GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_ROPE }; + std::initializer_list rms_norm_mul_rope_set_rows_ops = { GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_ROPE, GGML_OP_VIEW, GGML_OP_SET_ROWS }; + + if (is_equal(rms_norm_mul_rope_set_rows_ops, ops) && ggml_can_fuse_subgraph(cgraph, node_idx, ops, { node_idx + 4 })) { + const ggml_tensor * rms_norm = cgraph->nodes[node_idx]; + const ggml_tensor * mul = cgraph->nodes[node_idx + 1]; + const ggml_tensor * rope = cgraph->nodes[node_idx + 2]; + const ggml_tensor * view = cgraph->nodes[node_idx + 3]; + const ggml_tensor * set_rows = cgraph->nodes[node_idx + 4]; + + if (ggml_check_edges(cgraph, node_idx, {{1, 0, 0}, {2, 0, 1}, {3, 0, 2}, {4, 0, 3}}) && + ggml_cuda_should_fuse_rms_norm_mul_rope(rms_norm, mul, rope) && + ggml_cuda_should_fuse_rope_set_rows(rope, view, set_rows)) { + int out_nodes[] = { node_idx + 4 }; + return ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, (int)ops.size(), out_nodes, 1); + } + } + + if (is_equal(rms_norm_mul_rope_ops, ops) && ggml_can_fuse(cgraph, node_idx, ops)) { + const ggml_tensor * rms_norm = cgraph->nodes[node_idx]; + const ggml_tensor * mul = cgraph->nodes[node_idx + 1]; + const ggml_tensor * rope = cgraph->nodes[node_idx + 2]; + + if (ggml_cuda_should_fuse_rms_norm_mul_rope(rms_norm, mul, rope)) { + int out_nodes[] = { node_idx + 2 }; + return ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, (int)ops.size(), out_nodes, 1); + } + return false; + } + std::initializer_list rope_set_rows_ops = { GGML_OP_ROPE, GGML_OP_VIEW, GGML_OP_SET_ROWS }; if (is_equal(rope_set_rows_ops, ops) && ggml_can_fuse_subgraph(cgraph, node_idx, ops, { node_idx + 2 })) { @@ -2988,7 +3064,8 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, const ggml_tensor * set_rows = cgraph->nodes[node_idx + 2]; if (ggml_cuda_should_fuse_rope_set_rows(rope, view, set_rows)) { - return true; + int out_nodes[] = { node_idx + 2 }; + return ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, (int)ops.size(), out_nodes, 1); } } @@ -3840,6 +3917,16 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph return fused_node_count - 1; } + if (ggml_cuda_can_fuse(cgraph, i, { GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_ROPE, GGML_OP_VIEW, GGML_OP_SET_ROWS }, {})) { + ggml_cuda_op_rms_norm_mul_rope_fused(*cuda_ctx, node, cgraph->nodes[i + 1], cgraph->nodes[i + 2], cgraph->nodes[i + 4]); + return 4; + } + + if (ggml_cuda_can_fuse(cgraph, i, { GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_ROPE }, {})) { + ggml_cuda_op_rms_norm_mul_rope_fused(*cuda_ctx, node, cgraph->nodes[i + 1], cgraph->nodes[i + 2], nullptr); + return 2; + } + if (ggml_cuda_can_fuse(cgraph, i, { GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_ADD }, {})) { ggml_cuda_op_rms_norm_fused_add(*cuda_ctx, node, cgraph->nodes[i + 1], cgraph->nodes[i + 2]); return 2; diff --git a/ggml/src/ggml-cuda/rope.cu b/ggml/src/ggml-cuda/rope.cu index e20a5cb6b..504c6b818 100644 --- a/ggml/src/ggml-cuda/rope.cu +++ b/ggml/src/ggml-cuda/rope.cu @@ -670,3 +670,238 @@ void ggml_cuda_op_rope_back(ggml_backend_cuda_context & ctx, ggml_tensor * dst) void ggml_cuda_op_rope_fused(ggml_backend_cuda_context & ctx, ggml_tensor * rope, ggml_tensor * set_rows) { ggml_cuda_op_rope_impl(ctx, rope, set_rows); } + +// fused RMS_NORM + MUL + ROPE (+ VIEW + SET_ROWS) +// one block per row: block_reduce gives the norm scale, then each thread applies mul and rope to the elements it owns +template +static __global__ void rms_norm_mul_rope_f32( + const float * x, D * dst, const int ncols, + const int64_t s01, const int64_t s02, const int64_t s03, + const int64_t s1, const int64_t s2, const int64_t s3, + const float eps, + const float * mul, + const int64_t mul_s01, const int64_t mul_s02, const int64_t mul_s03, + const uint3 mul_ncols_packed, const uint3 mul_nrows_packed, + const uint3 mul_nchannels_packed, const uint3 mul_nsamples_packed, + const int n_dims, const int32_t * pos, + const float freq_scale, const float ext_factor, const float attn_factor, + const rope_corr_dims corr_dims, const float theta_scale, + const float * freq_factors, + const int64_t * row_indices, const int set_rows_stride, + const bool is_neox) { + ggml_cuda_pdl_lc(); + const int row = blockIdx.x; + const int channel = blockIdx.y; + const int sample = blockIdx.z; + const int tid = threadIdx.x; + + x += sample*s03 + channel*s02 + row*s01; + + const uint32_t mul_row = fastmodulo(row, mul_nrows_packed); + const uint32_t mul_channel = fastmodulo(channel, mul_nchannels_packed); + const uint32_t mul_sample = fastmodulo(sample, mul_nsamples_packed); + mul += mul_sample*mul_s03 + mul_channel*mul_s02 + mul_row*mul_s01; + + float tmp = 0.0f; + + ggml_cuda_pdl_sync(); + for (int col = tid; col < ncols; col += block_size) { + const float xi = x[col]; + tmp += xi * xi; + } + + extern __shared__ float s_sum[]; + tmp = block_reduce(tmp, s_sum); + + const float scale = rsqrtf(tmp/ncols + eps); + + int64_t idst = sample*s3 + channel*s2 + row*s1; + if (set_rows_stride != 0) { + idst = row*s1 + row_indices[channel]*set_rows_stride; + } + dst += idst; + + for (int i0 = 2*tid; i0 < ncols; i0 += 2*block_size) { + int ix0; + int ix1; + if (is_neox && i0 < n_dims) { + ix0 = i0/2; + ix1 = i0/2 + n_dims/2; + } else { + ix0 = i0 + 0; + ix1 = i0 + 1; + } + + const float x0 = scale * x[ix0] * mul[fastmodulo(ix0, mul_ncols_packed)]; + const float x1 = scale * x[ix1] * mul[fastmodulo(ix1, mul_ncols_packed)]; + + if (i0 >= n_dims) { + dst[ix0] = ggml_cuda_cast(x0); + dst[ix1] = ggml_cuda_cast(x1); + continue; + } + + const float theta_base = pos[channel]*powf(theta_scale, i0/2.0f); + const float freq_factor = has_ff ? freq_factors[i0/2] : 1.0f; + + float cos_theta; + float sin_theta; + rope_yarn(theta_base/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor, cos_theta, sin_theta); + + dst[ix0] = ggml_cuda_cast(x0*cos_theta - x1*sin_theta); + dst[ix1] = ggml_cuda_cast(x0*sin_theta + x1*cos_theta); + } +} + +template +static void rms_norm_mul_rope_cuda( + const float * x, D * dst, + const int ncols, const int nrows, const int nchannels, const int nsamples, + const int64_t s01, const int64_t s02, const int64_t s03, + const int64_t s1, const int64_t s2, const int64_t s3, + const float eps, + const float * mul, + const int64_t mul_s01, const int64_t mul_s02, const int64_t mul_s03, + const uint32_t mul_ncols, const uint32_t mul_nrows, + const uint32_t mul_nchannels, const uint32_t mul_nsamples, + const int n_dims, const int32_t * pos, + const float freq_scale, const float freq_base, const float ext_factor, const float attn_factor, + const rope_corr_dims corr_dims, + const float * freq_factors, + const int64_t * row_indices, const int set_rows_stride, + const bool is_neox, cudaStream_t stream) { + GGML_ASSERT(ncols % 2 == 0); + + const dim3 blocks_num(nrows, nchannels, nsamples); + + const float theta_scale = powf(freq_base, -2.0f/n_dims); + + const uint3 mul_ncols_packed = init_fastdiv_values(mul_ncols); + const uint3 mul_nrows_packed = init_fastdiv_values(mul_nrows); + const uint3 mul_nchannels_packed = init_fastdiv_values(mul_nchannels); + const uint3 mul_nsamples_packed = init_fastdiv_values(mul_nsamples); + + if (ncols < 1024) { + const dim3 block_dims(256, 1, 1); + const ggml_cuda_kernel_launch_params launch_params = {blocks_num, block_dims, 32*sizeof(float), stream}; + if (freq_factors == nullptr) { + ggml_cuda_kernel_launch(rms_norm_mul_rope_f32<256, false, D>, launch_params, + x, dst, ncols, s01, s02, s03, s1, s2, s3, eps, mul, mul_s01, mul_s02, mul_s03, + mul_ncols_packed, mul_nrows_packed, mul_nchannels_packed, mul_nsamples_packed, + n_dims, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale, + freq_factors, row_indices, set_rows_stride, is_neox); + } else { + ggml_cuda_kernel_launch(rms_norm_mul_rope_f32<256, true, D>, launch_params, + x, dst, ncols, s01, s02, s03, s1, s2, s3, eps, mul, mul_s01, mul_s02, mul_s03, + mul_ncols_packed, mul_nrows_packed, mul_nchannels_packed, mul_nsamples_packed, + n_dims, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale, + freq_factors, row_indices, set_rows_stride, is_neox); + } + } else { + const dim3 block_dims(1024, 1, 1); + const ggml_cuda_kernel_launch_params launch_params = {blocks_num, block_dims, 32*sizeof(float), stream}; + if (freq_factors == nullptr) { + ggml_cuda_kernel_launch(rms_norm_mul_rope_f32<1024, false, D>, launch_params, + x, dst, ncols, s01, s02, s03, s1, s2, s3, eps, mul, mul_s01, mul_s02, mul_s03, + mul_ncols_packed, mul_nrows_packed, mul_nchannels_packed, mul_nsamples_packed, + n_dims, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale, + freq_factors, row_indices, set_rows_stride, is_neox); + } else { + ggml_cuda_kernel_launch(rms_norm_mul_rope_f32<1024, true, D>, launch_params, + x, dst, ncols, s01, s02, s03, s1, s2, s3, eps, mul, mul_s01, mul_s02, mul_s03, + mul_ncols_packed, mul_nrows_packed, mul_nchannels_packed, mul_nsamples_packed, + n_dims, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale, + freq_factors, row_indices, set_rows_stride, is_neox); + } + } +} + +void ggml_cuda_op_rms_norm_mul_rope_fused(ggml_backend_cuda_context & ctx, + ggml_tensor * rms_norm, ggml_tensor * mul, ggml_tensor * rope, ggml_tensor * set_rows) { + const ggml_tensor * x = rms_norm->src[0]; + const ggml_tensor * mul_src = mul->src[0] == rms_norm ? mul->src[1] : mul->src[0]; + + float eps = 0.0f; + memcpy(&eps, rms_norm->op_params, sizeof(float)); + GGML_ASSERT(eps >= 0.0f); + + GGML_ASSERT(x->type == GGML_TYPE_F32); + GGML_ASSERT(mul_src->type == GGML_TYPE_F32); + GGML_ASSERT(rope->type == GGML_TYPE_F32); + + void * dst_d = rope->data; + ggml_type dst_type = rope->type; + const int64_t * row_indices = nullptr; + int set_rows_stride = 0; + + if (set_rows != nullptr) { + dst_d = set_rows->data; + dst_type = set_rows->type; + row_indices = (const int64_t *) set_rows->src[1]->data; + set_rows_stride = set_rows->nb[1] / ggml_type_size(set_rows->type); + } + + const int n_dims = ((const int32_t *) rope->op_params)[1]; + const int mode = ((const int32_t *) rope->op_params)[2]; + const int n_ctx_orig = ((const int32_t *) rope->op_params)[4]; + + float freq_base; + float freq_scale; + float ext_factor; + float attn_factor; + float beta_fast; + float beta_slow; + + memcpy(&freq_base, (const int32_t *) rope->op_params + 5, sizeof(float)); + memcpy(&freq_scale, (const int32_t *) rope->op_params + 6, sizeof(float)); + memcpy(&ext_factor, (const int32_t *) rope->op_params + 7, sizeof(float)); + memcpy(&attn_factor, (const int32_t *) rope->op_params + 8, sizeof(float)); + memcpy(&beta_fast, (const int32_t *) rope->op_params + 9, sizeof(float)); + memcpy(&beta_slow, (const int32_t *) rope->op_params + 10, sizeof(float)); + + const bool is_neox = mode & GGML_ROPE_TYPE_NEOX; + + const int32_t * pos = (const int32_t *) rope->src[1]->data; + + const float * freq_factors = rope->src[2] != nullptr ? (const float *) rope->src[2]->data : nullptr; + + rope_corr_dims corr_dims; + ggml_rope_yarn_corr_dims(n_dims, n_ctx_orig, freq_base, beta_fast, beta_slow, corr_dims.v); + + const size_t ts0 = ggml_type_size(x->type); + GGML_ASSERT(x->nb[0] == ts0); + const int64_t s01 = x->nb[1] / ts0; + const int64_t s02 = x->nb[2] / ts0; + const int64_t s03 = x->nb[3] / ts0; + + const size_t ts_mul = ggml_type_size(mul_src->type); + GGML_ASSERT(mul_src->nb[0] == ts_mul); + const int64_t mul_s01 = mul_src->nb[1] / ts_mul; + const int64_t mul_s02 = mul_src->nb[2] / ts_mul; + const int64_t mul_s03 = mul_src->nb[3] / ts_mul; + + const size_t ts_dst = ggml_type_size(rope->type); + const int64_t s1 = rope->nb[1] / ts_dst; + const int64_t s2 = rope->nb[2] / ts_dst; + const int64_t s3 = rope->nb[3] / ts_dst; + + cudaStream_t stream = ctx.stream(); + + if (dst_type == GGML_TYPE_F32) { + rms_norm_mul_rope_cuda((const float *) x->data, (float *) dst_d, + x->ne[0], x->ne[1], x->ne[2], x->ne[3], s01, s02, s03, s1, s2, s3, eps, + (const float *) mul_src->data, mul_s01, mul_s02, mul_s03, + mul_src->ne[0], mul_src->ne[1], mul_src->ne[2], mul_src->ne[3], + n_dims, pos, freq_scale, freq_base, ext_factor, attn_factor, corr_dims, + freq_factors, row_indices, set_rows_stride, is_neox, stream); + } else if (dst_type == GGML_TYPE_F16) { + rms_norm_mul_rope_cuda((const float *) x->data, (half *) dst_d, + x->ne[0], x->ne[1], x->ne[2], x->ne[3], s01, s02, s03, s1, s2, s3, eps, + (const float *) mul_src->data, mul_s01, mul_s02, mul_s03, + mul_src->ne[0], mul_src->ne[1], mul_src->ne[2], mul_src->ne[3], + n_dims, pos, freq_scale, freq_base, ext_factor, attn_factor, corr_dims, + freq_factors, row_indices, set_rows_stride, is_neox, stream); + } else { + GGML_ABORT("fatal error"); + } +} diff --git a/ggml/src/ggml-cuda/rope.cuh b/ggml/src/ggml-cuda/rope.cuh index 72af086cd..7ce2d71c5 100644 --- a/ggml/src/ggml-cuda/rope.cuh +++ b/ggml/src/ggml-cuda/rope.cuh @@ -7,3 +7,5 @@ void ggml_cuda_op_rope(ggml_backend_cuda_context & ctx, ggml_tensor * dst); void ggml_cuda_op_rope_back(ggml_backend_cuda_context & ctx, ggml_tensor * dst); void ggml_cuda_op_rope_fused(ggml_backend_cuda_context & ctx, ggml_tensor * dst, ggml_tensor * set_rows); + +void ggml_cuda_op_rms_norm_mul_rope_fused(ggml_backend_cuda_context & ctx, ggml_tensor * rms_norm, ggml_tensor * mul, ggml_tensor * rope, ggml_tensor * set_rows); diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 6688debd4..14a234060 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -2584,6 +2584,7 @@ struct test_rms_norm_mul_rope : public test_case { const float eps; const bool multi_add; // test a sequence of adds feeding into rms_norm const bool set_rows; + const bool broadcast; // multiply by a 1D [ne0] weight, as model norm weights are int mode; std::string op_desc(ggml_tensor * t) override { @@ -2594,12 +2595,12 @@ struct test_rms_norm_mul_rope : public test_case { bool run_whole_graph() override { return true; } std::string vars() override { - return VARS_TO_STR5(ne, eps, multi_add, set_rows, mode); + return VARS_TO_STR6(ne, eps, multi_add, set_rows, broadcast, mode); } test_rms_norm_mul_rope(std::array ne, float eps = 1e-6f, bool multi_add = false, - bool set_rows = false, int mode = GGML_ROPE_TYPE_NORMAL) - : ne(ne), eps(eps), multi_add(multi_add), set_rows(set_rows), mode(mode) {} + bool set_rows = false, bool broadcast = false, int mode = GGML_ROPE_TYPE_NORMAL) + : ne(ne), eps(eps), multi_add(multi_add), set_rows(set_rows), broadcast(broadcast), mode(mode) {} ggml_tensor * build_graph(ggml_context * ctx) override { ggml_tensor * a = ggml_new_tensor_4d(ctx, GGML_TYPE_F32, ne[0], ne[1], ne[2], 1); @@ -2610,7 +2611,9 @@ struct test_rms_norm_mul_rope : public test_case { a = ggml_add(ctx, ggml_add(ctx, a, b), c); } - a = ggml_mul(ctx, ggml_rms_norm(ctx, a, eps), b); + ggml_tensor * w = broadcast ? ggml_new_tensor_1d(ctx, GGML_TYPE_F32, ne[0]) : b; + + a = ggml_mul(ctx, ggml_rms_norm(ctx, a, eps), w); ggml_tensor * pos = ggml_new_tensor_1d(ctx, GGML_TYPE_I32, ne[2]); @@ -8756,16 +8759,18 @@ static std::vector> make_test_cases_eval() { for (auto multi_add : {false, true}) { for (auto set_rows : {false, true}) { - for (auto rope : {GGML_ROPE_TYPE_NORMAL, GGML_ROPE_TYPE_NEOX}) { - test_cases.emplace_back(new test_rms_norm_mul_rope({768, 1, 1, 1}, 1e-6f, multi_add, set_rows, rope)); - test_cases.emplace_back(new test_rms_norm_mul_rope({768, 3, 1, 1}, 1e-6f, multi_add, set_rows, rope)); - test_cases.emplace_back(new test_rms_norm_mul_rope({768, 3, 5, 1}, 1e-6f, multi_add, set_rows, rope)); - test_cases.emplace_back(new test_rms_norm_mul_rope({128, 32, 2, 1}, 1e-6f, multi_add, set_rows, rope)); - test_cases.emplace_back(new test_rms_norm_mul_rope({128, 4, 2, 1}, 1e-6f, multi_add, set_rows, rope)); - test_cases.emplace_back(new test_rms_norm_mul_rope({128, 32, 50, 1}, 1e-6f, multi_add, set_rows, rope)); - test_cases.emplace_back(new test_rms_norm_mul_rope({128, 4, 50, 1}, 1e-6f, multi_add, set_rows, rope)); - test_cases.emplace_back(new test_rms_norm_mul_rope({8192, 2, 2, 1}, 1e-6f, multi_add, set_rows, rope)); - test_cases.emplace_back(new test_rms_norm_mul_rope({8192, 2, 2, 1}, 1e-6f, multi_add, set_rows, rope)); + for (auto broadcast : {false, true}) { + for (auto rope : {GGML_ROPE_TYPE_NORMAL, GGML_ROPE_TYPE_NEOX}) { + test_cases.emplace_back(new test_rms_norm_mul_rope({768, 1, 1, 1}, 1e-6f, multi_add, set_rows, broadcast, rope)); + test_cases.emplace_back(new test_rms_norm_mul_rope({768, 3, 1, 1}, 1e-6f, multi_add, set_rows, broadcast, rope)); + test_cases.emplace_back(new test_rms_norm_mul_rope({768, 3, 5, 1}, 1e-6f, multi_add, set_rows, broadcast, rope)); + test_cases.emplace_back(new test_rms_norm_mul_rope({128, 32, 2, 1}, 1e-6f, multi_add, set_rows, broadcast, rope)); + test_cases.emplace_back(new test_rms_norm_mul_rope({128, 4, 2, 1}, 1e-6f, multi_add, set_rows, broadcast, rope)); + test_cases.emplace_back(new test_rms_norm_mul_rope({128, 32, 50, 1}, 1e-6f, multi_add, set_rows, broadcast, rope)); + test_cases.emplace_back(new test_rms_norm_mul_rope({128, 4, 50, 1}, 1e-6f, multi_add, set_rows, broadcast, rope)); + test_cases.emplace_back(new test_rms_norm_mul_rope({8192, 2, 2, 1}, 1e-6f, multi_add, set_rows, broadcast, rope)); + test_cases.emplace_back(new test_rms_norm_mul_rope({8192, 2, 2, 1}, 1e-6f, multi_add, set_rows, broadcast, rope)); + } } } } From 7ba604f1cb61cd14898138e9abc0b4ff2601f180 Mon Sep 17 00:00:00 2001 From: Pascal Date: Sun, 9 Aug 2026 00:42:50 +0200 Subject: [PATCH 05/16] server: report the isolate working directory from get_info (#26773) * server: report the isolate working directory from get_info Without an explicit cwd, get_info fell back to the server process working directory even when a tools runtime was configured. That named a host path no tool would ever run in, since an isolate starts in a directory of its own. It now asks the isolate for its working directory in that case, and keeps the process one only when the tools run on the host. * remove redundant comment --------- Co-authored-by: Xuan-Son Nguyen --- tools/server/server-tools.cpp | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/tools/server/server-tools.cpp b/tools/server/server-tools.cpp index d5e696434..27c663fdd 100644 --- a/tools/server/server-tools.cpp +++ b/tools/server/server-tools.cpp @@ -1669,8 +1669,13 @@ struct server_tool_get_info : server_tool { std::string cwd = json_value(params, "cwd", std::string()); if (cwd.empty()) { - std::error_code ec; - cwd = path_to_utf8(fs::current_path(ec)); + if (json_value(params, "runtime", std::string()).empty()) { + std::error_code ec; + cwd = path_to_utf8(fs::current_path(ec)); + } else { + auto pwd = io->run({"pwd"}, SERVER_TOOL_GET_INFO_MAX_OUTPUT, SERVER_TOOL_GET_INFO_TIMEOUT); + cwd = pwd.exit_code == 0 && !pwd.timed_out ? string_strip(pwd.output) : "unknown"; + } } return { From 61141f1487e63d9b22aec193131253bb2ea0800c Mon Sep 17 00:00:00 2001 From: Aaron Teo Date: Sun, 9 Aug 2026 18:15:28 +0800 Subject: [PATCH 06/16] ci: rm `GGML_HIP_ROCWMMA_FATTN` (#26760) Signed-off-by: Aaron Teo --- .devops/rocm.Dockerfile | 1 - .github/workflows/build-cuda-ubuntu.yml | 1 - .github/workflows/build-cuda-windows.yml | 1 - .github/workflows/release.yml | 2 -- ci/run.sh | 2 +- 5 files changed, 1 insertion(+), 6 deletions(-) diff --git a/.devops/rocm.Dockerfile b/.devops/rocm.Dockerfile index a8bc4e1fc..20f6ad636 100644 --- a/.devops/rocm.Dockerfile +++ b/.devops/rocm.Dockerfile @@ -57,7 +57,6 @@ COPY --from=web /app/tools/ui/dist tools/ui/dist RUN HIPCXX="$(hipconfig -l)/clang" HIP_PATH="$(hipconfig -R)" \ cmake -S . -B build \ -DGGML_HIP=ON \ - -DGGML_HIP_ROCWMMA_FATTN=ON \ -DAMDGPU_TARGETS="$ROCM_DOCKER_ARCH" \ -DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON \ -DCMAKE_BUILD_TYPE=Release -DLLAMA_BUILD_TESTS=OFF \ diff --git a/.github/workflows/build-cuda-ubuntu.yml b/.github/workflows/build-cuda-ubuntu.yml index 6271b22cb..2528b1857 100644 --- a/.github/workflows/build-cuda-ubuntu.yml +++ b/.github/workflows/build-cuda-ubuntu.yml @@ -99,7 +99,6 @@ jobs: run: | cmake -B build -S . \ -DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \ - -DGGML_HIP_ROCWMMA_FATTN=ON \ -DGPU_TARGETS="gfx1030" \ -DGGML_HIP=ON cmake --build build --config Release -j $(nproc) diff --git a/.github/workflows/build-cuda-windows.yml b/.github/workflows/build-cuda-windows.yml index e9e941421..367a3a854 100644 --- a/.github/workflows/build-cuda-windows.yml +++ b/.github/workflows/build-cuda-windows.yml @@ -150,7 +150,6 @@ jobs: -DLLAMA_BUILD_BORINGSSL=ON ` -DROCM_DIR="${env:HIP_PATH}" ` -DGGML_HIP=ON ` - -DGGML_HIP_ROCWMMA_FATTN=ON ` -DGPU_TARGETS="gfx1100" ` -DGGML_RPC=ON cmake --build build -j ${env:NUMBER_OF_PROCESSORS} diff --git a/.github/workflows/release.yml b/.github/workflows/release.yml index f587ed93c..968d2d4b7 100644 --- a/.github/workflows/release.yml +++ b/.github/workflows/release.yml @@ -1229,7 +1229,6 @@ jobs: -DGPU_TARGETS="${{ matrix.gpu_targets }}" \ -DGGML_HIP=ON \ -DHIP_PLATFORM=amd \ - -DGGML_HIP_ROCWMMA_FATTN=ON \ -DHF_UI_VERSION=${{ needs.get-version.outputs.ui_version }} \ ${{ env.CMAKE_ARGS }} cmake --build build --config Release -j $(nproc) @@ -1353,7 +1352,6 @@ jobs: -DGGML_NATIVE=OFF ` -DGGML_CPU=OFF ` -DGPU_TARGETS="${{ matrix.gpu_targets }}" ` - -DGGML_HIP_ROCWMMA_FATTN=ON ` -DGGML_HIP=ON ` -DHF_UI_VERSION=${{ needs.get-version.outputs.ui_version }} ` -DLLAMA_BUILD_BORINGSSL=ON diff --git a/ci/run.sh b/ci/run.sh index 8506bb408..f6c7eb0d5 100755 --- a/ci/run.sh +++ b/ci/run.sh @@ -92,7 +92,7 @@ if [ ! -z ${GG_BUILD_CUDA} ]; then fi if [ ! -z ${GG_BUILD_ROCM} ]; then - CMAKE_EXTRA="${CMAKE_EXTRA} -DCMAKE_HIP_COMPILER=$(hipconfig -l)/clang -DGGML_HIP=ON -DGGML_HIP_ROCWMMA_FATTN=ON" + CMAKE_EXTRA="${CMAKE_EXTRA} -DCMAKE_HIP_COMPILER=$(hipconfig -l)/clang -DGGML_HIP=ON" if [ -z ${GG_BUILD_AMDGPU_TARGETS} ]; then echo "Missing GG_BUILD_AMDGPU_TARGETS, please set it to your GPU architecture (e.g. gfx90a, gfx1100, etc.)" exit 1 From 08659901c43b51de735740f1cf61bb82fbe0c4e4 Mon Sep 17 00:00:00 2001 From: Hao-Chen2337 <2113996104@qq.com> Date: Sun, 9 Aug 2026 18:16:53 +0800 Subject: [PATCH 07/16] ggml-cpu : fix missing Q5_0 dispatch in SpaceMiT backend (#26792) --- ggml/src/ggml-cpu/spacemit/ime.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/ggml/src/ggml-cpu/spacemit/ime.cpp b/ggml/src/ggml-cpu/spacemit/ime.cpp index 9563ea3e4..29d683270 100644 --- a/ggml/src/ggml-cpu/spacemit/ime.cpp +++ b/ggml/src/ggml-cpu/spacemit/ime.cpp @@ -195,6 +195,7 @@ template class tensor_ case GGML_TYPE_Q4_K: case GGML_TYPE_Q6_K: case GGML_TYPE_Q8_0: + case GGML_TYPE_Q5_0: case GGML_TYPE_Q5_1: case GGML_TYPE_Q5_K: //case GGML_TYPE_MXFP4: @@ -214,6 +215,7 @@ template class tensor_ case GGML_TYPE_Q4_K: case GGML_TYPE_Q6_K: case GGML_TYPE_Q8_0: + case GGML_TYPE_Q5_0: case GGML_TYPE_Q5_1: case GGML_TYPE_Q5_K: //case GGML_TYPE_MXFP4: From 936918514ce522b553c0fd80b169a6440e6096c6 Mon Sep 17 00:00:00 2001 From: Xuan-Son Nguyen Date: Sun, 9 Aug 2026 16:51:21 +0200 Subject: [PATCH 08/16] ci: add pr-draft-label (#26801) --- .github/workflows/pr-draft-label.yml | 23 +++++++++++++++++++++++ 1 file changed, 23 insertions(+) create mode 100644 .github/workflows/pr-draft-label.yml diff --git a/.github/workflows/pr-draft-label.yml b/.github/workflows/pr-draft-label.yml new file mode 100644 index 000000000..d2594c823 --- /dev/null +++ b/.github/workflows/pr-draft-label.yml @@ -0,0 +1,23 @@ +name: Convert PR to draft + +on: + pull_request_target: + types: [labeled] + +permissions: + pull-requests: write + issues: write + contents: write # required for "gh pr ready" command, see https://github.com/cli/cli/issues/8910 + +jobs: + convert-to-draft: + if: github.event.label.name == 'draft' && github.event.pull_request.draft == false + runs-on: ubuntu-slim + steps: + - name: Convert PR to draft + env: + GH_TOKEN: ${{ secrets.GITHUB_TOKEN }} + PR_URL: ${{ github.event.pull_request.html_url }} + run: | + gh pr ready --undo "$PR_URL" + gh pr edit "$PR_URL" --remove-label draft From 74ce15741b420b8d6f12e720398458b576c51c2c Mon Sep 17 00:00:00 2001 From: Pascal Date: Sun, 9 Aug 2026 21:20:23 +0200 Subject: [PATCH 09/16] ui: degrade the working directory picker when file search is off (#26811) The picker mounts whenever a cwd-aware builtin tool is enabled, so it can open while file_glob_search is not served or was disabled by the user. Every typed query then fired a search that could only fail with a raw error. Gate the debounced search on the tool state, the same way the mention picker does, and show a message in place of the results list that explains why search is unavailable. Manual entry with Enter still commits a directory. The Browse button and the search scope footer are hidden as well: Browse resolves the picked folder name through file_glob_search, and the client-side toggle would not stop that call. --- .../ChatForm/ChatFormWorkingDirectory.svelte | 26 +++++++++++++++---- 1 file changed, 21 insertions(+), 5 deletions(-) diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormWorkingDirectory.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormWorkingDirectory.svelte index f8069c93e..108cf1b19 100644 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormWorkingDirectory.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormWorkingDirectory.svelte @@ -67,6 +67,20 @@ const pickerSupported = typeof window !== 'undefined' && typeof window.showDirectoryPicker === 'function'; + // When the server does not serve file_glob_search or the user disabled + // it, the picker still opens for manual entry but explains why search is + // unavailable instead of firing searches that would only fail. Browse is + // hidden too: it resolves the picked folder name through the same tool. + const fileSearchKey = $derived(toolsStore.getPermissionKey(BuiltInTool.FILE_GLOB_SEARCH)); + const fileSearchEnabled = $derived( + fileSearchKey !== null && toolsStore.isToolEnabled(fileSearchKey) + ); + const searchUnavailableMessage = $derived( + fileSearchKey === null + ? 'File search is unavailable on this server - type a full path and press Enter' + : 'File search is disabled - type a full path and press Enter, or enable "Search files" in Settings > Tools' + ); + let searchInputRef: HTMLInputElement | null = $state(null); let queryResults = $state([]); @@ -98,7 +112,7 @@ if (!isOpen) return; const q = query.trim(); nav.reset(-1); - if (q) { + if (q && fileSearchEnabled) { search.run(q); } else { search.cancel(); @@ -123,7 +137,7 @@ // children too, so path navigation does not require a trailing slash. const search = useDebouncedSearch({ debounceMs: SEARCH_DEBOUNCE_MS, - canRun: () => isOpen, + canRun: () => isOpen && fileSearchEnabled, getQuery: () => query.trim(), run: async (q, signal, isCurrent) => { const trimmed = q.trim(); @@ -340,7 +354,9 @@ class="w-full" /> - {#if query.trim() && (search.isSearching || queryResults.length > 0 || searchError)} + {#if !fileSearchEnabled} +
{searchUnavailableMessage}
+ {:else if query.trim() && (search.isSearching || queryResults.length > 0 || searchError)} {/if} - {#if pickerSupported} + {#if pickerSupported && fileSearchEnabled}