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 <thichthat@gmail.com>
CUDA: fuse rms_norm + mul + rope (+ view + set_rows) (#26767 )
2026-08-09 21:24:52 +02:00 · 2026-08-09 00:42:50 +02:00 · 2026-08-09 00:32:37 +08:00 · 2026-08-08 16:36:21 +02:00 · 2026-08-08 16:35:53 +02:00 · 2026-08-08 07:40:04 +03:00
22 changed files with 1027 additions and 127 deletions
@@ -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:<image>': spin up a new Docker container and reuse it for all invocations, clean up on server exit\n"
+        "  'docker-container:<id>': 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"
@@ -655,6 +655,7 @@ struct common_params {

    // enable built-in tools
    std::vector<std::string> 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
@@ -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_blck_f32_q8_0, QK8_0><<<num_blocks, 1, 0, stream>>>
+    cpy_f32_q<cpy_blck_f32_q8_0, QK8_0><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>
        (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_blck_q8_0_f32, QK8_0><<<num_blocks, 1, 0, stream>>>
+    cpy_q_f32<cpy_blck_q8_0_f32, QK8_0><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>
        (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_blck_f32_q4_0, QK4_0><<<num_blocks, 1, 0, stream>>>
+    cpy_f32_q<cpy_blck_f32_q4_0, QK4_0><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>
        (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<cpy_blck_q_f32<dequantize_q4_0, QK4_0>, QK4_0><<<num_blocks, 1, 0, stream>>>(
+    cpy_q_f32<cpy_blck_q_f32<dequantize_q4_0, QK4_0>, QK4_0><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>(
        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_blck_f32_q4_1, QK4_1><<<num_blocks, 1, 0, stream>>>
+    cpy_f32_q<cpy_blck_f32_q4_1, QK4_1><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>
        (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<cpy_blck_q_f32<dequantize_q4_1, QK4_1>, QK4_1><<<num_blocks, 1, 0, stream>>>(
+    cpy_q_f32<cpy_blck_q_f32<dequantize_q4_1, QK4_1>, QK4_1><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>(
        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_blck_f32_q5_0, QK5_0><<<num_blocks, 1, 0, stream>>>
+    cpy_f32_q<cpy_blck_f32_q5_0, QK5_0><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>
        (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<cpy_blck_q_f32<dequantize_q5_0, QK5_0>, QK5_0><<<num_blocks, 1, 0, stream>>>(
+    cpy_q_f32<cpy_blck_q_f32<dequantize_q5_0, QK5_0>, QK5_0><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>(
        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_blck_f32_q5_1, QK5_1><<<num_blocks, 1, 0, stream>>>
+    cpy_f32_q<cpy_blck_f32_q5_1, QK5_1><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>
        (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<cpy_blck_q_f32<dequantize_q5_1, QK5_1>, QK5_1><<<num_blocks, 1, 0, stream>>>(
+    cpy_q_f32<cpy_blck_q_f32<dequantize_q5_1, QK5_1>, QK5_1><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>(
        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_blck_f32_iq4_nl, QK4_NL><<<num_blocks, 1, 0, stream>>>
+    cpy_f32_q<cpy_blck_f32_iq4_nl, QK4_NL><<<num_blocks, CUDA_CPY_BLOCK_SIZE, 0, stream>>>
        (cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13);
 }

@@ -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<enum ggml_op> rms_norm_mul_rope_ops          = { GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_ROPE };
+    std::initializer_list<enum ggml_op> 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<enum ggml_op> 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;
@@ -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<true>(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 <int block_size, bool has_ff, typename D>
+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<block_reduce_method::SUM, block_size>(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<D>(x0);
+            dst[ix1] = ggml_cuda_cast<D>(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<true>(theta_base/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor, cos_theta, sin_theta);
+
+        dst[ix0] = ggml_cuda_cast<D>(x0*cos_theta - x1*sin_theta);
+        dst[ix1] = ggml_cuda_cast<D>(x0*sin_theta + x1*cos_theta);
+    }
+}
+
+template <typename D>
+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");
+    }
+}
@@ -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);
@@ -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<int64_t, 4> 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]);

@@ -8576,6 +8579,9 @@ static std::vector<std::unique_ptr<test_case>> 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}));
@@ -8753,16 +8759,18 @@ static std::vector<std::unique_ptr<test_case>> 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));
+                }
            }
        }
    }
@@ -54,7 +54,6 @@
 | `-ctv, --cache-type-v TYPE` | KV cache data type for V<br/>allowed values: f32, f16, bf16, q8_0, q4_0, q4_1, iq4_nl, q5_0, q5_1<br/>(default: f16)<br/>(env: LLAMA_ARG_CACHE_TYPE_V) |
 | `-dt, --defrag-thold N` | KV cache defragmentation threshold (DEPRECATED)<br/>(env: LLAMA_ARG_DEFRAG_THOLD) |
 | `-np, --parallel N` | number of parallel sequences to decode (default: 1)<br/>(env: LLAMA_ARG_N_PARALLEL) |
-| `--rpc SERVERS` | comma-separated list of RPC servers (host:port)<br/>(env: LLAMA_ARG_RPC) |
 | `--mlock` | DEPRECATED in favor of `--load-mode`: force system to keep model in RAM rather than swapping or compressing<br/>(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)<br/>(env: LLAMA_ARG_MMAP) |
 | `-dio, --direct-io, -ndio, --no-direct-io` | DEPRECATED in favor of `--load-mode`: use DirectIO if available<br/>(env: LLAMA_ARG_DIO) |
@@ -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<br/>allowed values: f32, f16, bf16, q8_0, q4_0, q4_1, iq4_nl, q5_0, q5_1<br/>(default: f16)<br/>(env: LLAMA_ARG_CACHE_TYPE_V) |
 | `-dt, --defrag-thold N` | KV cache defragmentation threshold (DEPRECATED)<br/>(env: LLAMA_ARG_DEFRAG_THOLD) |
 | `-np, --parallel N` | number of parallel sequences to decode (default: 1)<br/>(env: LLAMA_ARG_N_PARALLEL) |
-| `--rpc SERVERS` | comma-separated list of RPC servers (host:port)<br/>(env: LLAMA_ARG_RPC) |
 | `--mlock` | DEPRECATED in favor of `--load-mode`: force system to keep model in RAM rather than swapping or compressing<br/>(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)<br/>(env: LLAMA_ARG_MMAP) |
 | `-dio, --direct-io, -ndio, --no-direct-io` | DEPRECATED in favor of `--load-mode`: use DirectIO if available<br/>(env: LLAMA_ARG_DIO) |
@@ -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:<id>` 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):

@@ -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<br/>allowed values: f32, f16, bf16, q8_0, q4_0, q4_1, iq4_nl, q5_0, q5_1<br/>(default: f16)<br/>(env: LLAMA_ARG_CACHE_TYPE_K) |
 | `-ctv, --cache-type-v TYPE` | KV cache data type for V<br/>allowed values: f32, f16, bf16, q8_0, q4_0, q4_1, iq4_nl, q5_0, q5_1<br/>(default: f16)<br/>(env: LLAMA_ARG_CACHE_TYPE_V) |
 | `-dt, --defrag-thold N` | KV cache defragmentation threshold (DEPRECATED)<br/>(env: LLAMA_ARG_DEFRAG_THOLD) |
-| `--rpc SERVERS` | comma-separated list of RPC servers (host:port)<br/>(env: LLAMA_ARG_RPC) |
 | `--mlock` | DEPRECATED in favor of `--load-mode`: force system to keep model in RAM rather than swapping or compressing<br/>(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)<br/>(env: LLAMA_ARG_MMAP) |
 | `-dio, --direct-io, -ndio, --no-direct-io` | DEPRECATED in favor of `--load-mode`: use DirectIO if available<br/>(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)<br/>(env: LLAMA_ARG_UI_CONFIG) |
 | `--ui-config-file, --webui-config-file PATH` | JSON file that provides default UI settings (overrides UI defaults)<br/>(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)<br/>(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)<br/>specify "all" to enable all tools<br/>available tools: read_file, file_glob_search, grep_search, exec_shell_command, write_file, edit_file, get_datetime<br/>note: for security reasons, this will limit --cors-origins to localhost by default<br/>(env: LLAMA_ARG_TOOLS) |
+| `--tools-runtime OPTION` | experimental: run tools in a separate runtime environment (default: none, use host environment)<br/>available options:<br/>  'docker:<image>': spin up a new Docker container and reuse it for all invocations, clean up on server exit<br/>  'docker-container:<id>': use an existing Docker container by ID, won't stop on server exit<br/><br/>(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)<br/>specify "all" to enable all tools<br/>available tools: read_file, file_glob_search, grep_search, exec_shell_command, write_file, edit_file, get_datetime, get_info<br/>note: for security reasons, this will limit --cors-origins to localhost by default<br/>(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)<br/>note: for security reasons, this will limit --cors-origins to localhost by default<br/>(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)<br/>note: for security reasons, this will limit --cors-origins to localhost by default<br/>(env: LLAMA_ARG_MCP_SERVERS_JSON) |
@@ -10,12 +10,15 @@
 #include <ctime>
 #include <atomic>
 #include <cstring>
+#include <cstdint>
 #include <cstdlib>
 #include <algorithm>
+#include <iterator>
 #include <unordered_set>
 #include <tuple>
 #include <functional>
 #include <memory>
+#include <mutex>

 #if defined(_WIN32)
 #   ifndef NOMINMAX
@@ -71,6 +74,7 @@ json server_tool::to_json() const {
        {"permissions", json{
            {"write", permission_write}
        }},
+        {"uses_cwd", uses_cwd},
        {"definition", get_definition()},
    };
 }
@@ -127,6 +131,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 +176,85 @@ public:
            const std::function<bool(const std::string &)> & 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<std::string> & args,
+        size_t max_output,
+        int timeout_secs,
+        const std::function<bool(const std::string &)> & 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<bool> done{false};
+    std::atomic<bool> 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 +366,7 @@ public:
            size_t max_output,
            int timeout_secs,
            const std::function<bool(const std::string &)> & 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<bool> done{false};
-        std::atomic<bool> 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 +409,8 @@ private:
    }

    static const std::unordered_set<std::string> & junk_dir_names() {
-        static const std::unordered_set<std::string> names = {
-            ".git", ".svn", ".hg", "node_modules", "__pycache__",
-            ".venv", "venv", "dist", "build", "target", ".cache", ".idea", ".vscode",
-        };
+        static const std::unordered_set<std::string> names(
+            std::begin(SERVER_TOOL_JUNK_DIR_NAMES), std::end(SERVER_TOOL_JUNK_DIR_NAMES));
        return names;
    }

@@ -450,9 +473,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<uint64_t> tmp_counter{0};
+        fs::path tmp = tmp_dir / string_format(
+            "llama-tools-io-isolate-%zu-%llu.tmp",
+            std::hash<std::thread::id>{}(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<std::string> & args,
+            size_t max_output,
+            int timeout_secs,
+            const std::function<bool(const std::string &)> & on_chunk = nullptr) const override {
+        std::vector<std::string> 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<std::string> build_argv(const std::vector<std::string> & 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<std::string> & 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<std::string> with_cwd(const std::vector<std::string> & inner) const {
+        if (cwd.empty()) {
+            return inner;
+        }
+        // 127 is what a shell reports for a command it could not run
+        std::vector<std::string> 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<std::string> & 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<std::string> & 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<std::string> split_lines(const std::string & text, bool strip_dot_slash) {
+        std::vector<std::string> 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<std::string> 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<std::string> build_argv(const std::vector<std::string> & inner, bool needs_stdin) const override {
+        std::vector<std::string> 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<tools_io> make_tools_io(const json & params) {
-    std::string cwd = json_value(params, "cwd", std::string());
-    return std::make_unique<tools_io_basic>(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<tools_io_basic>(cwd);
+    }
+    if (runtime.rfind(SERVER_TOOL_RUNTIME_DOCKER_CONTAINER, 0) == 0) {
+        return std::make_unique<tools_io_docker>(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
@@ -476,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;
    }

@@ -564,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;
    }

@@ -678,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;
    }

@@ -830,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;
    }
@@ -861,8 +1153,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<std::string> args = {"cmd", "/c", command};
+        std::vector<std::string> args = !json_value(params, "runtime", std::string()).empty()
+            ? std::vector<std::string>{"sh", "-c", command}
+            : std::vector<std::string>{"cmd", "/c", command};
 #else
        std::vector<std::string> args = {"sh", "-c", command};
 #endif
@@ -905,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;
    }

@@ -947,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;
    }

@@ -1335,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;
    }

@@ -1355,19 +1653,29 @@ 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<std::string> args = !json_value(params, "runtime", std::string()).empty()
+            ? std::vector<std::string>{"uname", "-a"}
+            : std::vector<std::string>{"cmd", "/c", "ver"};
 #else
-        auto res = io->run({"uname", "-a"}, SERVER_TOOL_GET_INFO_MAX_OUTPUT, SERVER_TOOL_GET_INFO_TIMEOUT);
+        std::vector<std::string> 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";

        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 {
@@ -1461,6 +1769,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:<image> 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:<id> 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<std::mutex> 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 <image> 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::thread::id>{}(std::this_thread::get_id()));
+        fs::remove(cidfile, ec);
+
+        std::vector<std::string> 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<std::unique_ptr<server_tool>> & tools, const std::string & name, bool require_stream) {
    for (auto & t : tools) {
        if (t->name == name) {
@@ -1506,8 +1911,16 @@ static std::string get_header(const std::map<std::string, std::string> & headers
    return default_value;
 }

+server_tools::server_tools() = default;
+server_tools::~server_tools() = default;
+
 void server_tools::setup(const std::vector<std::string> & enabled_tools,
-                         server_mcp & mcp_mgr) {
+                         server_mcp & mcp_mgr,
+                         const std::string & tools_runtime) {
+    if (!tools_runtime.empty()) {
+        docker_runtime = std::make_unique<server_tools_docker_runtime>(tools_runtime);
+    }
+
    if (!enabled_tools.empty()) {
        if (!common_subproc::is_supported()) {
            throw std::runtime_error("subprocess is not enabled on this build");
@@ -1590,11 +2003,26 @@ void server_tools::setup(const std::vector<std::string> & 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:<id>";
+            // 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) {
@@ -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;
@@ -30,6 +31,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<std::unique_ptr<server_tool>> tools;

@@ -37,9 +40,16 @@ struct server_tools {
    server_response queue_res;
    std::atomic<int> res_id{0};

+    // set when --tools-runtime is configured; owns the docker container used to run tools, if any
+    std::unique_ptr<server_tools_docker_runtime> docker_runtime;
+
    void setup(const std::vector<std::string> & 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();
 };
@@ -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)");
        }
@@ -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()
@@ -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:
@@ -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 @@

 	<ContextGaugePopup />

-	{#if toolsStore.builtinTools.length > 0}
+	{#if toolsStore.hasEnabledCwdTools}
 		<ChatFormWorkingDirectory
 			directory={cwd}
 			isOpen={pickers.isWorkingDirectoryPickerOpen}
@@ -8,7 +8,7 @@ interface ChatCommandsOptions {
 	/** Gates `/prompt`. */
 	hasPrompts: () => 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',
@@ -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
 		})
 	);

@@ -27,6 +27,9 @@ class ToolsStore {
 	private _loading = $state(false);
 	private _error = $state<string | null>(null);
 	private _disabledTools = $state(new SvelteSet<string>());
+	// builtin tools that resolve their paths against the working directory,
+	// as declared by the server in its `/tools` listing
+	private _cwdAwareTools = $state(new SvelteSet<string>());
 	private _toolsEndpointUnreachable = $state(false);
 	private _serverHome = $state<string | null | undefined>(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<void> {
 		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;
@@ -292,6 +292,7 @@ export interface ServerBuiltinToolInfo {
 	permissions: {
 		write: boolean;
 	};
+	uses_cwd: boolean;
 	definition: OpenAIToolDefinition;
 }

@@ -21,7 +21,7 @@
 		focusInput: () => {},
 		getShowModelSelector: () => true,
 		hasPrompts: () => true,
-		hasBuiltinTools: () => true,
+		hasCwdTools: () => true,
 		getCwd: () => null,
 		getServerHome: () => null,
 		openModelSelector: () => {
Author	SHA1	Message	Date
Pascal GitHub Xuan-Son Nguyen	7ba604f1cb	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 <thichthat@gmail.com>	2026-08-09 00:42:50 +02:00
Rafail GiavrimisandGitHub	687e778927	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	2026-08-09 00:32:37 +08:00
PascalandGitHub	18f7ad7fc9	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.	2026-08-08 16:36:21 +02:00
Xuan-Son Nguyen GitHub Pascal	dd2c7c4471	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 <admin@serveurperso.com>	2026-08-08 16:35:53 +02:00
Rafail GiavrimisandGitHub	69bf643791	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	2026-08-08 07:40:04 +03:00