Merge commit '66001722aa' into concedo_experimental

# Conflicts:
#	README.md
#	docs/ops.md
#	docs/ops/SYCL.csv
#	examples/sycl/start-svr.sh
#	ggml/src/ggml-hexagon/ggml-hexagon.cpp
#	ggml/src/ggml-hexagon/htp/CMakeLists.txt
#	ggml/src/ggml-hexagon/htp/htp-ctx.h
#	ggml/src/ggml-hexagon/htp/htp-ops.h
#	ggml/src/ggml-hexagon/htp/main.c
#	ggml/src/ggml-hexagon/htp/unary-ops.c
#	ggml/src/ggml-opencl/CMakeLists.txt
#	ggml/src/ggml-opencl/ggml-opencl.cpp
#	ggml/src/ggml-opencl/kernels/cvt.cl
#	ggml/src/ggml-sycl/gated_delta_net.hpp
#	ggml/src/ggml-sycl/ggml-sycl.cpp
#	ggml/src/ggml-sycl/pad.cpp
#	ggml/src/ggml-sycl/ssm_conv.cpp
#	tests/test-backend-ops.cpp
#	tests/test-reasoning-budget.cpp
#	tools/server/README.md
#	tools/server/webui/src/lib/constants/settings-config.ts
This commit is contained in:
Concedo 2026-05-11 15:40:10 +08:00
commit 9b0b36b5ef
97 changed files with 8192 additions and 4853 deletions

View file

@ -369,9 +369,7 @@ common_peg_parser analyze_tools::build_tool_parser_tag_tagged(parser_build_conte
arguments.name_suffix) +
arguments.value_prefix +
(schema_info.resolves_to_string(param_schema) ?
p.tool_arg_string_value(p.schema(until_suffix,
"tool-" + name + "-arg-" + param_name + "-schema",
param_schema, true)) :
p.tool_arg_string_value(until_suffix) :
p.tool_arg_json_value(p.schema(
p.json(), "tool-" + name + "-arg-" + param_name + "-schema", param_schema, false)) +
p.space()) +

View file

@ -91,7 +91,7 @@ json common_chat_msg::to_json_oaicompat(bool concat_typed_text) const {
if (!content.empty()) {
jmsg["content"] = content;
} else if (!content_parts.empty()) {
if (concat_typed_text) {
if (concat_typed_text || contains_media()) {
std::string text;
bool last_was_media_marker = false;
// join parts with newline, do not add newline before or after media markers

View file

@ -94,6 +94,15 @@ struct common_chat_msg {
tool_name.empty() && tool_call_id.empty();
}
bool contains_media() const {
for (const auto & part : content_parts) {
if (part.type == "media_marker") {
return true;
}
}
return false;
}
void set_tool_call_ids(std::vector<std::string> & ids_cache,
const std::function<std::string()> & gen_tool_call_id) {
for (auto i = 0u; i < tool_calls.size(); i++) {

View file

@ -158,8 +158,6 @@ static void common_reasoning_budget_apply(struct llama_sampler * smpl, llama_tok
for (size_t i = 0; i < cur_p->size; i++) {
if (cur_p->data[i].id != forced) {
cur_p->data[i].logit = -INFINITY;
} else {
cur_p->data[i].logit = +INFINITY; // force the token
}
}
}

View file

@ -710,7 +710,7 @@ class ModelBase:
self._repack_nvfp4(name, weight, scale, scale2, input_scale)
# Flush any remaining experts (fallback if n_experts was unknown)
for bid, proj_type in expert_blocks.keys():
for bid, proj_type in list(expert_blocks.keys()):
self._flush_nvfp4_experts((bid, proj_type), expert_blocks, expert_scales, expert_input_scales, expert_shapes, bid, proj_type)
# Remove consumed tensors so get_tensors/modify_tensors won't see them
@ -718,7 +718,7 @@ class ModelBase:
self.model_tensors.pop(name, None)
# Remove any remaining unused auxiliary tensors
for name in self.model_tensors.keys():
for name in list(self.model_tensors.keys()):
if name.endswith((".k_scale", ".v_scale")):
del self.model_tensors[name]
@ -7988,13 +7988,37 @@ class Gemma4Model(Gemma3Model):
rope_freqs_full = torch.tensor(values, dtype=torch.float32)
yield (self.format_tensor_name(gguf.MODEL_TENSOR.ROPE_FREQS), rope_freqs_full)
def _generate_nvfp4_tensors(self):
# Gemma-4 stores a per-layer router.per_expert_scale ([n_expert]) that scales
# each expert's contribution. It's mathematically equivalent to a per-expert
# scalar on the down_proj output, which is exactly where ffn_down_exps_s is
# applied at inference. Fold it into each expert's NVFP4 weight_scale_2 so the
# existing NVFP4 path produces the right scales.
n_experts = self.find_hparam(["num_local_experts", "num_experts"], optional=True) or 0
for name in [n for n in self.model_tensors if n.endswith(".router.per_expert_scale")]:
bid_match = re.search(r"\.layers\.(\d+)\.", name)
if bid_match is None:
continue
bid = bid_match.group(1)
prefix = name[: name.index(f".layers.{bid}.") + len(f".layers.{bid}.")]
w2_targets = [f"{prefix}experts.{e}.down_proj.weight_scale_2" for e in range(n_experts)]
present = [w2 in self.model_tensors for w2 in w2_targets]
if not any(present):
continue
assert all(present), f"layer {bid}: partial NVFP4 quantization across experts"
r = self.model_tensors.pop(name)
for e, w2 in enumerate(w2_targets):
s = self.model_tensors[w2]
self.model_tensors[w2] = lambda s=s, r=r, i=e: s() * r()[i]
super()._generate_nvfp4_tensors()
@classmethod
def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
name, gen = item
if name.endswith("per_dim_scale") or name.endswith("layer_scalar"):
name = name + ".weight"
if ".experts." in name and not name.endswith(".weight"):
if ".experts." in name and not name.endswith((".weight", ".weight_scale", ".weight_scale_2", ".input_scale")):
name += ".weight"
return super().filter_tensors((name, gen))
@ -13684,6 +13708,27 @@ class DotsOCRVisionModel(MmprojModel):
yield from super().modify_tensors(data_torch, name, bid)
@ModelBase.register("Sarashina2VisionForCausalLM")
class Sarashina2VLTextModel(LlamaModel):
model_arch = gguf.MODEL_ARCH.LLAMA
@classmethod
def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
name, gen = item
if name.startswith("llm."):
name = name.replace("llm.", "", 1)
elif name.startswith("norm."):
return None
return super().filter_tensors((name, gen))
@ModelBase.register("Sarashina2VisionForCausalLM")
class Sarashina2VLVisionModel(Qwen2VLVisionModel):
def __init__(self, *args, **kwargs):
super().__init__(*args, **kwargs)
self.global_config['model_type'] = "qwen2_vl"
###### CONVERSION LOGIC ######
@ -13940,7 +13985,7 @@ def get_model_architecture(hparams: dict[str, Any], model_type: ModelType) -> st
# Step3-VL keeps text config under text_config but uses a custom top-level architecture.
# For text conversion we route to a dedicated text-only class.
# TODO: refactor this later to avoid adding exception here
if model_type == ModelType.TEXT and arch == "StepVLForConditionalGeneration":
if model_type == ModelType.TEXT and arch in ("StepVLForConditionalGeneration", "Sarashina2VisionForCausalLM"):
return arch
# if "architectures" is found in the sub-config, use that instead

127
docs/multi-gpu.md Normal file
View file

@ -0,0 +1,127 @@
# Using multiple GPUs with llama.cpp
This guide explains how to run [llama.cpp](https://github.com/ggml-org/llama.cpp) across more than one GPU. It covers the split modes, the command-line flags that control them, the limitations you need to know about, and ready-to-use recipes for `llama-cli` and `llama-server`.
The CLI arguments listed here are the same for both tools - or most llama.cpp binaries for that matter.
---
## When you need multi-GPU
Reach for multi-GPU when one of these is true:
- **The model doesn't fit in a single GPU's VRAM.** By spreading the weights across two or more GPUs the whole model can stay on accelerators. Otherwise part of the model will need to be run off of the comparatively slower system RAM.
- **You want more throughput.** By distributing the computation across multiple GPUs, each individual GPU has to do less work. This can result in better prefill and/or token generation performance, depending on the split mode and interconnect speed vs. the speed of an individual GPU.
---
## The split modes
Set with `--split-mode` / `-sm`.
| Mode | What it does | When to use |
|---|---|---|
| `none` | Use a single GPU only. Pick which one with `--main-gpu`. | You explicitly want to confine the model to one GPU even though more are visible. |
| `layer` (**default**) | Pipeline parallelism. Each GPU holds a contiguous slice of layers. The KV cache for layer *l* lives on the GPU that owns layer *l*. | Default and most compatible multi-GPU choice. You want more memory than a single GPU provides and your priority is a fast prefill. Can tolerate slow interconnect speeds between GPUs. |
| `row` | **Deprecated.** Older row-split tensor-parallel path with comparatively poor performance. Splits only dense weights across GPUs. Superseded by `tensor` which should be universally superior if it can be used. | Avoid in new deployments. |
| `tensor` | **EXPERIMENTAL.** Tensor parallelism that splits both weights *and* KV across the participating GPUs via a "meta device" abstraction. | You want more memory than a single GPU provides and your priority is fast token generation. Prefill speeds approach pipeline parallel speeds for large, dense models and fast GPU interconnect speeds. Treat as experimental as the code is less mature than pipeline parallelism. Performance should be good for multiple NVIDIA GPUs using the CUDA backend, no guarantees otherwise. |
> Pipeline parallel (`layer`) vs. tensor parallel (`tensor`): pipeline-parallel runs different layers on different GPUs and processes tokens sequentially through the pipeline. This minimizes data transfers between GPUs but requires many tokens to scale well. Tensor-parallel splits each layer across GPUs and does multiple cross-GPU reductions per layer. This enables parallelizing any workload but is much more bottlenecked by the GPU interconnect speed. Pipeline-parallel maximizes batch throughput; tensor-parallel minimizes latency.
---
## Command-line arguments reference
| Short | Long | Value | Default | Notes |
|---|---|---|---|---|
| `-sm` | `--split-mode` | `none` \| `layer` \| `tensor` | `layer` | See modes above. |
| `-ts` | `--tensor-split` | comma-separated proportions, e.g. `3,1` | mode-dependent | How much of the model goes to each GPU. If omitted, `layer`/`row` use automatic splitting proportional to memory, while `tensor` splits tensor segments evenly. With `3,1` on two GPUs, GPU 0 gets 75 %, GPU 1 gets 25 %. The values follow the order in `--device`. |
| `-mg` | `--main-gpu` | integer device index | `0` | The single GPU used in `--split-mode none`. |
| `-ngl` | `--n-gpu-layers` / `--gpu-layers` | integer \| `auto` \| `all` | `auto` | Maximum number of layers to keep in VRAM. Use `999` or `all` to push everything possible to the GPUs. |
| `-dev` | `--device` | comma-separated device names, or `none` | auto | Restrict which devices llama.cpp may use. See `--list-devices` for names. |
| | `--list-devices` | - | - | Print the available devices and their memory. Run this first to learn the names you'd pass to `--device`. |
| `-fa` | `--flash-attn` | `on` \| `off` \| `auto` | `auto` | Required when using `--split-mode tensor` and/or quantized V cache. Supported (and therefore enabled by default) for most combinations of models and backends. |
| `-ctk` | `--cache-type-k` | `f32` \| `f16` \| `bf16` \| `q8_0` \| `q4_0` \| ... | `f16` | KV cache type for K. |
| `-ctv` | `--cache-type-v` | same as `-ctk` | `f16` | KV cache type for V. |
| `-fit` | `--fit` | `on` \| `off` | `on` | Auto-fit unset args to device memory. **Not supported with `tensor`. You may need to manually set the `--ctx-size` to make the model fit.** |
As for any CUDA program, the environment variable `CUDA_VISIBLE_DEVICES` can be used to control which GPUs to use for the CUDA backend: if you set it, llama.cpp only sees the specified GPUs. Use `--device` for selecting GPUs from among those visible to llama.cpp, this works for any backend.
---
## Recipes
### 1. Default - pipeline parallel across all visible GPUs
```bash
llama-cli -m model.gguf
llama-server -m model.gguf
```
Easiest configuration. KV cache spreads across the GPUs along with the layers. `--fit` (on by default) sizes things automatically.
### 2. Pipeline parallel with a custom split ratio
```bash
llama-cli -m model.gguf -ts 3,1
```
Useful when GPUs have different memory: GPU 0 (3 parts) and GPU 1 (1 part). Proportions are normalized so `-ts 3,1` is the same as e.g. `-ts 75,25`.
### 3. Single-GPU mode, picking a specific GPU
```bash
llama-cli --list-devices
llama-cli -m model.gguf -dev CUDA1
```
Use only the device listed as `CUDA1` when calling with `--list-devices`.
### 4. Tensor parallelism (experimental)
```bash
llama-cli -m model.gguf -sm tensor -ctk f16 -ctv f16
```
- `--flash-attn off` or (`--flash-attn auto` resolving to `off` when it isn't supported) is a hard error.
- KV cache types must be non-quantized: `f32`, `f16`, or `bf16`. Support for quantized KV cache is not implemented and trying to use it will result in an error.
- Mark this configuration as experimental in your tooling: validate output quality before deploying.
- `--split-mode tensor`is not implemented for all architectures. The following will fail with *"LLAMA_SPLIT_MODE_TENSOR not implemented for architecture '...'"*:
- **MoE / hybrid:** Grok, MPT, OLMoE, DeepSeek2, GLM-DSA, Nemotron-H, Nemotron-H-MoE, Granite-Hybrid, LFM2-MoE, Minimax-M2, Mistral4, Kimi-Linear, Jamba, Falcon-H1
- **State-space / RWKV-style:** Mamba, Mamba2 (and the hybrid Mamba-attention models above)
- **Other:** PLAMO2, MiniCPM3, Gemma-3n, OLMo2, BitNet, T5
### 5. With NCCL
There's no runtime flag for NCCL - it's selected at build time (`-DGGML_CUDA_NCCL=ON`, this is the default). Note that NCCL is **not** automatically distributed with CUDA and you may need to install it manually - when in doubt check the CMake log to see whether or not it can find the package. When llama.cpp is compiled with NCCL support it uses it automatically for cross-GPU reductions in `tensor` mode. When NCCL is missing on a multi-GPU build, you'll see this one-time warning and performance will be lower:
```
NVIDIA Collective Communications Library (NCCL) is unavailable, multi GPU performance will be suboptimal
```
When using the "ROCm" backend (which is the ggml CUDA code translated for AMD via HIP), the AMD equivalent RCCL can be used by compiling with `-DGGML_HIP_RCCL=ON`. Note that RCCL is by default *disabled* because (unlike NCCL) it was not universally beneficial during testing.
### 6. With CUDA peer-to-peer access (`GGML_CUDA_P2P`)
CUDA peer-to-peer (P2P) lets GPUs transfer data directly between each other instead of going through system memory, which generally improves multi-GPU performance. It is **opt-in** at runtime - set the environment variable `GGML_CUDA_P2P` to any value to enable it:
```bash
GGML_CUDA_P2P=1 llama-cli -m model.gguf -sm tensor
```
P2P requires driver support (usually restricted to workstation/datacenter GPUs) and **may cause crashes or corrupted outputs on some motherboards or BIOS configurations** (e.g. when IOMMU is enabled). If you see instability after enabling it, unset the variable.
---
## Troubleshooting
| Symptom | How to fix |
|---|---|
| Startup error *"SPLIT_MODE_TENSOR requires flash_attn to be enabled"* | Add `-fa on` or remove `-fa off`. |
| Startup error *"simultaneous use of SPLIT_MODE_TENSOR and KV cache quantization not implemented"* | Use `-ctk f16 -ctv f16` (or `bf16`/`f32`) with `--split-mode tensor`. |
| Startup error *"LLAMA_SPLIT_MODE_TENSOR not implemented for architecture 'X'"* | Architecture not on the TENSOR allow-list. Use `--split-mode layer`. |
| Warning *"NCCL is unavailable, multi GPU performance will be suboptimal"* | llama.cpp wasn't built with NCCL. Either accept the lower performance or install NCCL and rebuild. |
| CUDA OOM at startup or during prefill in `--split-mode tensor` | Auto-fit is disabled in this mode, so reduce memory pressure yourself. In order from least to most disruptive: lower `--ctx-size` (`-c`) (KV cache is roughly proportional to `n_ctx`); for `llama-server`, lower `--parallel` (`-np`) (a slot KV cache is allocated per concurrent sequence); as a last resort, reduce `--n-gpu-layers` (`-ngl`) (the remaining layers run on CPU and inference will be much slower). |
| Performance is worse with multi-GPU than single-GPU | The performance is bottlenecked by GPU interconnect speed. For `--split-mode tensor`, verify that NCCL is being used. Try `--split-mode layer` (less communication than `tensor`). Increase GPU interconnect speed via more PCIe lanes or e.g. NVLink (if available). |
| GPU not used at all | `--n-gpu-layers` is `0` or too low - try explicitly setting `-ngl all`. Or you are accidentally hiding the GPUs via an environment variable like `CUDA_VISIBLE_DEVICES=-1`. Or your build doesn't include support for the relevant backend. |
| Crashes or corrupted outputs after setting `GGML_CUDA_P2P=1` | Some motherboards and BIOS settings (e.g. with IOMMU enabled) don't support CUDA peer-to-peer reliably. Unset `GGML_CUDA_P2P`. |

View file

@ -169,7 +169,7 @@ extern "C" {
// device type
enum ggml_backend_dev_type type;
// device id
// for PCI devices, this should be the PCI bus id formatted as "domain:bus:device.function" (e.g. "0000:01:00.0")
// for PCI devices, this should be the lower-case PCI bus id formatted as "domain:bus:device.function" (e.g. "0000:c1:00.0")
// if the id is unknown, this should be NULL
const char * device_id;
// device capabilities

View file

@ -972,7 +972,7 @@ static void ggml_backend_sched_print_assignments(ggml_backend_sched_t sched, str
}
if (sched->debug > 1) {
ggml_backend_t tensor_backend = ggml_backend_sched_get_tensor_backend(sched, node);
GGML_LOG_DEBUG("node #%3d (%10.10s): %20.20s (%5.5s) [%5.5s %8.8s] use=%d,c=%d:", i, ggml_op_name(node->op), node->name,
GGML_LOG_DEBUG("node #%3d (%10.10s): %20.20s (%5.5s) [%5.5s %8.8s] use=%d,c=%d:", i, ggml_op_desc(node), node->name,
fmt_size(ggml_nbytes(node)), tensor_backend ? ggml_backend_name(tensor_backend) : "NULL", GET_CAUSE(node),
graph->use_counts[ggml_hash_find(&graph->visited_hash_set, node)], node->flags & GGML_TENSOR_FLAG_COMPUTE ? 1 : 0);
for (int j = 0; j < GGML_MAX_SRC; j++) {

View file

@ -41,6 +41,7 @@ bool g_mul_mat_q = true;
#include "ggml-cuda/rope.cuh"
#include "ggml-cuda/roll.cuh"
#include "ggml-cuda/scale.cuh"
#include "ggml-cuda/snake.cuh"
#include "ggml-cuda/softcap.cuh"
#include "ggml-cuda/softmax.cuh"
#include "ggml-cuda/ssm-conv.cuh"
@ -3773,6 +3774,35 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
return 2;
}
// Snake activation: y = x + sin(a*x)^2 * inv_b
// Naive 5-op decomposition emitted by frontends: mul -> sin -> sqr -> mul -> add
if (ggml_can_fuse_subgraph(cgraph, i,
{ GGML_OP_MUL, GGML_OP_SIN, GGML_OP_SQR, GGML_OP_MUL, GGML_OP_ADD },
{ i + 4 })) {
const ggml_tensor * mul0 = cgraph->nodes[i];
const ggml_tensor * sqr = cgraph->nodes[i + 2];
const ggml_tensor * mul1 = cgraph->nodes[i + 3];
ggml_tensor * add = cgraph->nodes[i + 4];
// x carries the full activation shape, a is the broadcast operand
const ggml_tensor * x = ggml_are_same_shape(mul0, mul0->src[0]) ? mul0->src[0] : mul0->src[1];
const ggml_tensor * a = (x == mul0->src[0]) ? mul0->src[1] : mul0->src[0];
// mul1 reads sqr and inv_b in either operand order
const ggml_tensor * inv_b = (mul1->src[0] == sqr) ? mul1->src[1] : mul1->src[0];
// closure check: the trailing add must read the same x as the leading mul
const ggml_tensor * x_in_add = (add->src[0] == mul1) ? add->src[1] : add->src[0];
const bool type_ok = (x->type == GGML_TYPE_F32 || x->type == GGML_TYPE_F16 || x->type == GGML_TYPE_BF16);
const bool shape_ok = ggml_are_same_shape(a, inv_b) && a->ne[0] == 1 && a->ne[1] == x->ne[1];
if (type_ok && shape_ok && x_in_add == x && add->type == x->type) {
ggml_cuda_op_snake_fused(*cuda_ctx, x, a, inv_b, add);
return 4;
}
}
// multi-(add or mul)
if (node->op == GGML_OP_ADD || node->op == GGML_OP_MUL) {
int n_fuse = 0;
@ -5456,6 +5486,9 @@ ggml_backend_reg_t ggml_backend_cuda_reg() {
char pci_bus_id[32] = {};
CUDA_CHECK(cudaDeviceGetPCIBusId(pci_bus_id, sizeof(pci_bus_id), i));
dev_ctx->pci_bus_id = pci_bus_id;
for (char & c : dev_ctx->pci_bus_id) {
c = std::tolower(c);
}
dev_ctx->op_offload_min_batch_size = min_batch_size;
ggml_backend_dev_t dev = new ggml_backend_device {

View file

@ -54,15 +54,31 @@ void ggml_cuda_out_prod(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
const int64_t dps2 = ne2 / ne02;
const int64_t dps3 = ne3 / ne03;
// TODO batched matrix multiplication
for (int64_t i3 = 0; i3 < ne3; ++i3) {
for (int64_t i2 = 0; i2 < ne2; ++i2) {
if (dps2 == 1 && ne2 > 1) {
// src0 has uniform stride s02 along dim 2; batch the inner loop with a strided GEMM
GGML_ASSERT(ne2 <= std::numeric_limits<int>::max());
const int batch_count = (int) ne2;
for (int64_t i3 = 0; i3 < ne3; ++i3) {
CUBLAS_CHECK(
cublasSgemm(handle, CUBLAS_OP_N, src1_cublas_op,
cublasSgemmStridedBatched(handle, CUBLAS_OP_N, src1_cublas_op,
ne0, ne1, ne01,
&alpha, src0_d + (i3/dps3)*s03 + (i2/dps2)*s02, lda,
src1_d + i3 *s13 + i2 *s12, ldb,
&beta, dst_d + i3 *s3 + i2 *s2, ldc));
&alpha, src0_d + (i3/dps3)*s03, lda, s02,
src1_d + i3 *s13, ldb, s12,
&beta, dst_d + i3 *s3, ldc, s2,
batch_count));
}
} else {
// Fallback: ne2 == 1 (no batching benefit) or dps2 > 1 (src0 broadcast along dim 2
// with non-uniform stride; would need cublasSgemmBatched with pointer arrays).
for (int64_t i3 = 0; i3 < ne3; ++i3) {
for (int64_t i2 = 0; i2 < ne2; ++i2) {
CUBLAS_CHECK(
cublasSgemm(handle, CUBLAS_OP_N, src1_cublas_op,
ne0, ne1, ne01,
&alpha, src0_d + (i3/dps3)*s03 + (i2/dps2)*s02, lda,
src1_d + i3 *s13 + i2 *s12, ldb,
&beta, dst_d + i3 *s3 + i2 *s2, ldc));
}
}
}
}

View file

@ -0,0 +1,72 @@
#include "snake.cuh"
#include "convert.cuh"
// Fused Snake activation: y = x + sin^2(a * x) * inv_b
// x: [T, C] (T contiguous), a: [1, C], inv_b: [1, C]
// Supports F32, F16, BF16 data with F32 compute.
template <typename T>
static __global__ void snake_kernel(
const T * __restrict__ x,
const float * __restrict__ a,
const float * __restrict__ inv_b,
T * __restrict__ dst,
const int total,
const uint3 T_len_fastdiv) {
const int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx >= total) return;
const int c = (int) fastdiv((uint32_t) idx, T_len_fastdiv);
const float xi = ggml_cuda_cast<float>(x[idx]);
const float s = sinf(a[c] * xi);
dst[idx] = ggml_cuda_cast<T>(xi + s * s * inv_b[c]);
}
// Internal launcher with explicit x/a/inv_b/dst tensors.
// Shared by the public op (reads dst->src) and the fusion path (explicit args).
static void launch_snake(ggml_backend_cuda_context & ctx,
const ggml_tensor * x,
const ggml_tensor * a,
const ggml_tensor * inv_b,
ggml_tensor * dst) {
const float * a_d = (const float *)a->data;
const float * inv_b_d = (const float *)inv_b->data;
const int T = (int)x->ne[0];
const int C = (int)x->ne[1];
const int total = T * C;
const uint3 T_len_fastdiv = init_fastdiv_values((uint64_t) T);
const int block_size = 256;
const int grid_size = (total + block_size - 1) / block_size;
cudaStream_t stream = ctx.stream();
switch (x->type) {
case GGML_TYPE_F32: {
snake_kernel<<<grid_size, block_size, 0, stream>>>(
(const float *)x->data, a_d, inv_b_d, (float *)dst->data, total, T_len_fastdiv);
} break;
case GGML_TYPE_F16: {
snake_kernel<<<grid_size, block_size, 0, stream>>>(
(const half *)x->data, a_d, inv_b_d, (half *)dst->data, total, T_len_fastdiv);
} break;
case GGML_TYPE_BF16: {
snake_kernel<<<grid_size, block_size, 0, stream>>>(
(const nv_bfloat16 *)x->data, a_d, inv_b_d, (nv_bfloat16 *)dst->data, total, T_len_fastdiv);
} break;
default:
GGML_ABORT("snake: unsupported type");
}
}
// Fusion entry: caller supplies x/a/inv_b explicitly from the matched
// mul -> sin -> sqr -> mul -> add pattern. The dst is the trailing add output.
void ggml_cuda_op_snake_fused(ggml_backend_cuda_context & ctx,
const ggml_tensor * x,
const ggml_tensor * a,
const ggml_tensor * inv_b,
ggml_tensor * dst) {
launch_snake(ctx, x, a, inv_b, dst);
}

View file

@ -0,0 +1,8 @@
#include "common.cuh"
// Fusion entry point. Caller supplies x/a/inv_b explicitly.
void ggml_cuda_op_snake_fused(ggml_backend_cuda_context & ctx,
const ggml_tensor * x,
const ggml_tensor * a,
const ggml_tensor * inv_b,
ggml_tensor * dst);

View file

@ -48,6 +48,7 @@
#define cublasSetMathMode(handle, mode) CUBLAS_STATUS_SUCCESS
#define cublasSetStream hipblasSetStream
#define cublasSgemm hipblasSgemm
#define cublasSgemmStridedBatched hipblasSgemmStridedBatched
#define cublasStatus_t hipblasStatus_t
#define cublasOperation_t hipblasOperation_t
#define cudaDevAttrCooperativeLaunch hipDeviceAttributeCooperativeLaunch

View file

@ -32,6 +32,7 @@
#define cublasSetMathMode mublasSetMathMode
#define cublasSetStream mublasSetStream
#define cublasSgemm mublasSgemm
#define cublasSgemmStridedBatched mublasSgemmStridedBatched
#define cublasStatus_t mublasStatus_t
#define cublasOperation_t mublasOperation_t
#define cublasGetStatusString mublasGetStatusString

View file

@ -0,0 +1,955 @@
#include <math.h>
#include <stdint.h>
#include <string.h>
#include "hvx-utils.h"
#define GGML_COMMON_DECL_C
#include "ggml-common.h"
#include "htp-ctx.h"
#ifndef MIN
#define MIN(a, b) ((a) < (b) ? (a) : (b))
#endif
#define HTP_GDN_MAX_SV 128
struct htp_gdn_context {
struct htp_ops_context * octx;
uint32_t rows_per_thread;
size_t state_bytes;
bool use_vtcm;
uint8_t * vtcm_state_base;
size_t vtcm_state_per_thread;
};
static inline float gdn_mul_dot_f32(float * restrict dst, const float * restrict mul,
const float * restrict dot, uint32_t n) {
HVX_Vector acc = Q6_V_vzero();
const uint32_t epv = 128 / sizeof(float);
const uint32_t nvec = n / epv;
const uint32_t tail = n % epv;
for (uint32_t i = 0; i < nvec; ++i) {
HVX_Vector vd = hvx_vmemu(dst + i * epv);
HVX_Vector vm = hvx_vmem(mul + i * epv);
HVX_Vector vdot = hvx_vmem(dot + i * epv);
HVX_Vector out = hvx_vec_mul_f32_f32(vd, vm);
hvx_vmemu(dst + i * epv) = out;
acc = hvx_vec_add_f32_f32(acc, hvx_vec_mul_f32_f32(out, vdot));
}
if (tail) {
const uint32_t off = nvec * epv;
HVX_Vector vd = hvx_vmemu(dst + off);
HVX_Vector vm = hvx_vmem(mul + off);
HVX_Vector vdot = hvx_vmem(dot + off);
HVX_Vector out = hvx_vec_mul_f32_f32(vd, vm);
hvx_vec_store_u(dst + off, tail * sizeof(float), out);
HVX_VectorPred mask = Q6_Q_vsetq2_R(tail * sizeof(float));
HVX_Vector prod = hvx_vec_mul_f32_f32(out, vdot);
acc = hvx_vec_add_f32_f32(acc, Q6_V_vmux_QVV(mask, prod, Q6_V_vzero()));
}
return hvx_vec_get_f32(hvx_vec_reduce_sum_f32(acc));
}
static inline float gdn_mul_scalar_dot_f32(float * restrict dst, float mul,
const float * restrict dot, uint32_t n) {
HVX_Vector acc = Q6_V_vzero();
const HVX_Vector vmul = hvx_vec_splat_f32(mul);
const uint32_t epv = 128 / sizeof(float);
const uint32_t nvec = n / epv;
const uint32_t tail = n % epv;
for (uint32_t i = 0; i < nvec; ++i) {
HVX_Vector vd = hvx_vmemu(dst + i * epv);
HVX_Vector vdot = hvx_vmem(dot + i * epv);
HVX_Vector out = hvx_vec_mul_f32_f32(vd, vmul);
hvx_vmemu(dst + i * epv) = out;
acc = hvx_vec_add_f32_f32(acc, hvx_vec_mul_f32_f32(out, vdot));
}
if (tail) {
const uint32_t off = nvec * epv;
HVX_Vector vd = hvx_vmemu(dst + off);
HVX_Vector vdot = hvx_vmem(dot + off);
HVX_Vector out = hvx_vec_mul_f32_f32(vd, vmul);
hvx_vec_store_u(dst + off, tail * sizeof(float), out);
HVX_VectorPred mask = Q6_Q_vsetq2_R(tail * sizeof(float));
HVX_Vector prod = hvx_vec_mul_f32_f32(out, vdot);
acc = hvx_vec_add_f32_f32(acc, Q6_V_vmux_QVV(mask, prod, Q6_V_vzero()));
}
return hvx_vec_get_f32(hvx_vec_reduce_sum_f32(acc));
}
static inline float gdn_add_scaled_dot_f32(float * restrict dst, const float * restrict src,
float scale, const float * restrict dot, uint32_t n) {
HVX_Vector acc = Q6_V_vzero();
const HVX_Vector vscale = hvx_vec_splat_f32(scale);
const uint32_t epv = 128 / sizeof(float);
const uint32_t nvec = n / epv;
const uint32_t tail = n % epv;
for (uint32_t i = 0; i < nvec; ++i) {
HVX_Vector vd = hvx_vmemu(dst + i * epv);
HVX_Vector vs = hvx_vmem(src + i * epv);
HVX_Vector vdot = hvx_vmem(dot + i * epv);
HVX_Vector out = hvx_vec_add_f32_f32(vd, hvx_vec_mul_f32_f32(vs, vscale));
hvx_vmemu(dst + i * epv) = out;
acc = hvx_vec_add_f32_f32(acc, hvx_vec_mul_f32_f32(out, vdot));
}
if (tail) {
const uint32_t off = nvec * epv;
HVX_Vector vd = hvx_vmemu(dst + off);
HVX_Vector vs = hvx_vmem(src + off);
HVX_Vector vdot = hvx_vmem(dot + off);
HVX_Vector out = hvx_vec_add_f32_f32(vd, hvx_vec_mul_f32_f32(vs, vscale));
hvx_vec_store_u(dst + off, tail * sizeof(float), out);
HVX_VectorPred mask = Q6_Q_vsetq2_R(tail * sizeof(float));
HVX_Vector prod = hvx_vec_mul_f32_f32(out, vdot);
acc = hvx_vec_add_f32_f32(acc, Q6_V_vmux_QVV(mask, prod, Q6_V_vzero()));
}
return hvx_vec_get_f32(hvx_vec_reduce_sum_f32(acc));
}
static inline void gdn_mul_dot4_f32(float * restrict dst0, float * restrict dst1,
float * restrict dst2, float * restrict dst3, const float * restrict mul,
const float * restrict dot, uint32_t n, float * restrict sums) {
HVX_Vector acc0 = Q6_V_vzero();
HVX_Vector acc1 = Q6_V_vzero();
HVX_Vector acc2 = Q6_V_vzero();
HVX_Vector acc3 = Q6_V_vzero();
const uint32_t epv = 128 / sizeof(float);
const uint32_t nvec = n / epv;
const uint32_t tail = n % epv;
for (uint32_t i = 0; i < nvec; ++i) {
HVX_Vector vm = hvx_vmem(mul + i * epv);
HVX_Vector vdot = hvx_vmem(dot + i * epv);
HVX_Vector out0 = hvx_vec_mul_f32_f32(hvx_vmemu(dst0 + i * epv), vm);
HVX_Vector out1 = hvx_vec_mul_f32_f32(hvx_vmemu(dst1 + i * epv), vm);
HVX_Vector out2 = hvx_vec_mul_f32_f32(hvx_vmemu(dst2 + i * epv), vm);
HVX_Vector out3 = hvx_vec_mul_f32_f32(hvx_vmemu(dst3 + i * epv), vm);
hvx_vmemu(dst0 + i * epv) = out0;
hvx_vmemu(dst1 + i * epv) = out1;
hvx_vmemu(dst2 + i * epv) = out2;
hvx_vmemu(dst3 + i * epv) = out3;
acc0 = hvx_vec_add_f32_f32(acc0, hvx_vec_mul_f32_f32(out0, vdot));
acc1 = hvx_vec_add_f32_f32(acc1, hvx_vec_mul_f32_f32(out1, vdot));
acc2 = hvx_vec_add_f32_f32(acc2, hvx_vec_mul_f32_f32(out2, vdot));
acc3 = hvx_vec_add_f32_f32(acc3, hvx_vec_mul_f32_f32(out3, vdot));
}
if (tail) {
const uint32_t off = nvec * epv;
HVX_Vector vm = hvx_vmem(mul + off);
HVX_Vector vdot = hvx_vmem(dot + off);
HVX_VectorPred mask = Q6_Q_vsetq2_R(tail * sizeof(float));
HVX_Vector zero = Q6_V_vzero();
HVX_Vector out0 = hvx_vec_mul_f32_f32(hvx_vmemu(dst0 + off), vm);
HVX_Vector out1 = hvx_vec_mul_f32_f32(hvx_vmemu(dst1 + off), vm);
HVX_Vector out2 = hvx_vec_mul_f32_f32(hvx_vmemu(dst2 + off), vm);
HVX_Vector out3 = hvx_vec_mul_f32_f32(hvx_vmemu(dst3 + off), vm);
hvx_vec_store_u(dst0 + off, tail * sizeof(float), out0);
hvx_vec_store_u(dst1 + off, tail * sizeof(float), out1);
hvx_vec_store_u(dst2 + off, tail * sizeof(float), out2);
hvx_vec_store_u(dst3 + off, tail * sizeof(float), out3);
acc0 = hvx_vec_add_f32_f32(acc0, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out0, vdot), zero));
acc1 = hvx_vec_add_f32_f32(acc1, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out1, vdot), zero));
acc2 = hvx_vec_add_f32_f32(acc2, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out2, vdot), zero));
acc3 = hvx_vec_add_f32_f32(acc3, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out3, vdot), zero));
}
HVX_Vector_x4 acc = { .v = { acc0, acc1, acc2, acc3 } };
hvx_vec_store_u(sums, 4 * sizeof(float), hvx_vec_reduce_sum_f32x4(acc));
}
static inline void gdn_mul_scalar_dot4_f32(float * restrict dst0, float * restrict dst1,
float * restrict dst2, float * restrict dst3, float mul,
const float * restrict dot, uint32_t n, float * restrict sums) {
HVX_Vector acc0 = Q6_V_vzero();
HVX_Vector acc1 = Q6_V_vzero();
HVX_Vector acc2 = Q6_V_vzero();
HVX_Vector acc3 = Q6_V_vzero();
const HVX_Vector vmul = hvx_vec_splat_f32(mul);
const uint32_t epv = 128 / sizeof(float);
const uint32_t nvec = n / epv;
const uint32_t tail = n % epv;
for (uint32_t i = 0; i < nvec; ++i) {
HVX_Vector vdot = hvx_vmem(dot + i * epv);
HVX_Vector out0 = hvx_vec_mul_f32_f32(hvx_vmemu(dst0 + i * epv), vmul);
HVX_Vector out1 = hvx_vec_mul_f32_f32(hvx_vmemu(dst1 + i * epv), vmul);
HVX_Vector out2 = hvx_vec_mul_f32_f32(hvx_vmemu(dst2 + i * epv), vmul);
HVX_Vector out3 = hvx_vec_mul_f32_f32(hvx_vmemu(dst3 + i * epv), vmul);
hvx_vmemu(dst0 + i * epv) = out0;
hvx_vmemu(dst1 + i * epv) = out1;
hvx_vmemu(dst2 + i * epv) = out2;
hvx_vmemu(dst3 + i * epv) = out3;
acc0 = hvx_vec_add_f32_f32(acc0, hvx_vec_mul_f32_f32(out0, vdot));
acc1 = hvx_vec_add_f32_f32(acc1, hvx_vec_mul_f32_f32(out1, vdot));
acc2 = hvx_vec_add_f32_f32(acc2, hvx_vec_mul_f32_f32(out2, vdot));
acc3 = hvx_vec_add_f32_f32(acc3, hvx_vec_mul_f32_f32(out3, vdot));
}
if (tail) {
const uint32_t off = nvec * epv;
HVX_Vector vdot = hvx_vmem(dot + off);
HVX_VectorPred mask = Q6_Q_vsetq2_R(tail * sizeof(float));
HVX_Vector zero = Q6_V_vzero();
HVX_Vector out0 = hvx_vec_mul_f32_f32(hvx_vmemu(dst0 + off), vmul);
HVX_Vector out1 = hvx_vec_mul_f32_f32(hvx_vmemu(dst1 + off), vmul);
HVX_Vector out2 = hvx_vec_mul_f32_f32(hvx_vmemu(dst2 + off), vmul);
HVX_Vector out3 = hvx_vec_mul_f32_f32(hvx_vmemu(dst3 + off), vmul);
hvx_vec_store_u(dst0 + off, tail * sizeof(float), out0);
hvx_vec_store_u(dst1 + off, tail * sizeof(float), out1);
hvx_vec_store_u(dst2 + off, tail * sizeof(float), out2);
hvx_vec_store_u(dst3 + off, tail * sizeof(float), out3);
acc0 = hvx_vec_add_f32_f32(acc0, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out0, vdot), zero));
acc1 = hvx_vec_add_f32_f32(acc1, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out1, vdot), zero));
acc2 = hvx_vec_add_f32_f32(acc2, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out2, vdot), zero));
acc3 = hvx_vec_add_f32_f32(acc3, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out3, vdot), zero));
}
HVX_Vector_x4 acc = { .v = { acc0, acc1, acc2, acc3 } };
hvx_vec_store_u(sums, 4 * sizeof(float), hvx_vec_reduce_sum_f32x4(acc));
}
static inline void gdn_add_scaled_dot4_f32(float * restrict dst0, float * restrict dst1,
float * restrict dst2, float * restrict dst3, const float * restrict src,
const float * restrict scale, const float * restrict dot, uint32_t n,
float * restrict sums) {
HVX_Vector acc0 = Q6_V_vzero();
HVX_Vector acc1 = Q6_V_vzero();
HVX_Vector acc2 = Q6_V_vzero();
HVX_Vector acc3 = Q6_V_vzero();
const HVX_Vector scale0 = hvx_vec_splat_f32(scale[0]);
const HVX_Vector scale1 = hvx_vec_splat_f32(scale[1]);
const HVX_Vector scale2 = hvx_vec_splat_f32(scale[2]);
const HVX_Vector scale3 = hvx_vec_splat_f32(scale[3]);
const uint32_t epv = 128 / sizeof(float);
const uint32_t nvec = n / epv;
const uint32_t tail = n % epv;
for (uint32_t i = 0; i < nvec; ++i) {
HVX_Vector vs = hvx_vmem(src + i * epv);
HVX_Vector vdot = hvx_vmem(dot + i * epv);
HVX_Vector out0 = hvx_vec_add_f32_f32(hvx_vmemu(dst0 + i * epv), hvx_vec_mul_f32_f32(vs, scale0));
HVX_Vector out1 = hvx_vec_add_f32_f32(hvx_vmemu(dst1 + i * epv), hvx_vec_mul_f32_f32(vs, scale1));
HVX_Vector out2 = hvx_vec_add_f32_f32(hvx_vmemu(dst2 + i * epv), hvx_vec_mul_f32_f32(vs, scale2));
HVX_Vector out3 = hvx_vec_add_f32_f32(hvx_vmemu(dst3 + i * epv), hvx_vec_mul_f32_f32(vs, scale3));
hvx_vmemu(dst0 + i * epv) = out0;
hvx_vmemu(dst1 + i * epv) = out1;
hvx_vmemu(dst2 + i * epv) = out2;
hvx_vmemu(dst3 + i * epv) = out3;
acc0 = hvx_vec_add_f32_f32(acc0, hvx_vec_mul_f32_f32(out0, vdot));
acc1 = hvx_vec_add_f32_f32(acc1, hvx_vec_mul_f32_f32(out1, vdot));
acc2 = hvx_vec_add_f32_f32(acc2, hvx_vec_mul_f32_f32(out2, vdot));
acc3 = hvx_vec_add_f32_f32(acc3, hvx_vec_mul_f32_f32(out3, vdot));
}
if (tail) {
const uint32_t off = nvec * epv;
HVX_Vector vs = hvx_vmem(src + off);
HVX_Vector vdot = hvx_vmem(dot + off);
HVX_VectorPred mask = Q6_Q_vsetq2_R(tail * sizeof(float));
HVX_Vector zero = Q6_V_vzero();
HVX_Vector out0 = hvx_vec_add_f32_f32(hvx_vmemu(dst0 + off), hvx_vec_mul_f32_f32(vs, scale0));
HVX_Vector out1 = hvx_vec_add_f32_f32(hvx_vmemu(dst1 + off), hvx_vec_mul_f32_f32(vs, scale1));
HVX_Vector out2 = hvx_vec_add_f32_f32(hvx_vmemu(dst2 + off), hvx_vec_mul_f32_f32(vs, scale2));
HVX_Vector out3 = hvx_vec_add_f32_f32(hvx_vmemu(dst3 + off), hvx_vec_mul_f32_f32(vs, scale3));
hvx_vec_store_u(dst0 + off, tail * sizeof(float), out0);
hvx_vec_store_u(dst1 + off, tail * sizeof(float), out1);
hvx_vec_store_u(dst2 + off, tail * sizeof(float), out2);
hvx_vec_store_u(dst3 + off, tail * sizeof(float), out3);
acc0 = hvx_vec_add_f32_f32(acc0, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out0, vdot), zero));
acc1 = hvx_vec_add_f32_f32(acc1, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out1, vdot), zero));
acc2 = hvx_vec_add_f32_f32(acc2, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out2, vdot), zero));
acc3 = hvx_vec_add_f32_f32(acc3, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out3, vdot), zero));
}
HVX_Vector_x4 acc = { .v = { acc0, acc1, acc2, acc3 } };
hvx_vec_store_u(sums, 4 * sizeof(float), hvx_vec_reduce_sum_f32x4(acc));
}
static inline void gdn_mul_dot8_f32(float * restrict dst0, float * restrict dst1,
float * restrict dst2, float * restrict dst3, float * restrict dst4,
float * restrict dst5, float * restrict dst6, float * restrict dst7,
const float * restrict mul, const float * restrict dot, uint32_t n,
float * restrict sums) {
HVX_Vector acc0 = Q6_V_vzero();
HVX_Vector acc1 = Q6_V_vzero();
HVX_Vector acc2 = Q6_V_vzero();
HVX_Vector acc3 = Q6_V_vzero();
HVX_Vector acc4 = Q6_V_vzero();
HVX_Vector acc5 = Q6_V_vzero();
HVX_Vector acc6 = Q6_V_vzero();
HVX_Vector acc7 = Q6_V_vzero();
const uint32_t epv = 128 / sizeof(float);
const uint32_t nvec = n / epv;
const uint32_t tail = n % epv;
for (uint32_t i = 0; i < nvec; ++i) {
HVX_Vector vm = hvx_vmem(mul + i * epv);
HVX_Vector vdot = hvx_vmem(dot + i * epv);
HVX_Vector out0 = hvx_vec_mul_f32_f32(hvx_vmemu(dst0 + i * epv), vm);
HVX_Vector out1 = hvx_vec_mul_f32_f32(hvx_vmemu(dst1 + i * epv), vm);
HVX_Vector out2 = hvx_vec_mul_f32_f32(hvx_vmemu(dst2 + i * epv), vm);
HVX_Vector out3 = hvx_vec_mul_f32_f32(hvx_vmemu(dst3 + i * epv), vm);
HVX_Vector out4 = hvx_vec_mul_f32_f32(hvx_vmemu(dst4 + i * epv), vm);
HVX_Vector out5 = hvx_vec_mul_f32_f32(hvx_vmemu(dst5 + i * epv), vm);
HVX_Vector out6 = hvx_vec_mul_f32_f32(hvx_vmemu(dst6 + i * epv), vm);
HVX_Vector out7 = hvx_vec_mul_f32_f32(hvx_vmemu(dst7 + i * epv), vm);
hvx_vmemu(dst0 + i * epv) = out0;
hvx_vmemu(dst1 + i * epv) = out1;
hvx_vmemu(dst2 + i * epv) = out2;
hvx_vmemu(dst3 + i * epv) = out3;
hvx_vmemu(dst4 + i * epv) = out4;
hvx_vmemu(dst5 + i * epv) = out5;
hvx_vmemu(dst6 + i * epv) = out6;
hvx_vmemu(dst7 + i * epv) = out7;
acc0 = hvx_vec_add_f32_f32(acc0, hvx_vec_mul_f32_f32(out0, vdot));
acc1 = hvx_vec_add_f32_f32(acc1, hvx_vec_mul_f32_f32(out1, vdot));
acc2 = hvx_vec_add_f32_f32(acc2, hvx_vec_mul_f32_f32(out2, vdot));
acc3 = hvx_vec_add_f32_f32(acc3, hvx_vec_mul_f32_f32(out3, vdot));
acc4 = hvx_vec_add_f32_f32(acc4, hvx_vec_mul_f32_f32(out4, vdot));
acc5 = hvx_vec_add_f32_f32(acc5, hvx_vec_mul_f32_f32(out5, vdot));
acc6 = hvx_vec_add_f32_f32(acc6, hvx_vec_mul_f32_f32(out6, vdot));
acc7 = hvx_vec_add_f32_f32(acc7, hvx_vec_mul_f32_f32(out7, vdot));
}
if (tail) {
const uint32_t off = nvec * epv;
HVX_Vector vm = hvx_vmem(mul + off);
HVX_Vector vdot = hvx_vmem(dot + off);
HVX_VectorPred mask = Q6_Q_vsetq2_R(tail * sizeof(float));
HVX_Vector zero = Q6_V_vzero();
HVX_Vector out0 = hvx_vec_mul_f32_f32(hvx_vmemu(dst0 + off), vm);
HVX_Vector out1 = hvx_vec_mul_f32_f32(hvx_vmemu(dst1 + off), vm);
HVX_Vector out2 = hvx_vec_mul_f32_f32(hvx_vmemu(dst2 + off), vm);
HVX_Vector out3 = hvx_vec_mul_f32_f32(hvx_vmemu(dst3 + off), vm);
HVX_Vector out4 = hvx_vec_mul_f32_f32(hvx_vmemu(dst4 + off), vm);
HVX_Vector out5 = hvx_vec_mul_f32_f32(hvx_vmemu(dst5 + off), vm);
HVX_Vector out6 = hvx_vec_mul_f32_f32(hvx_vmemu(dst6 + off), vm);
HVX_Vector out7 = hvx_vec_mul_f32_f32(hvx_vmemu(dst7 + off), vm);
hvx_vec_store_u(dst0 + off, tail * sizeof(float), out0);
hvx_vec_store_u(dst1 + off, tail * sizeof(float), out1);
hvx_vec_store_u(dst2 + off, tail * sizeof(float), out2);
hvx_vec_store_u(dst3 + off, tail * sizeof(float), out3);
hvx_vec_store_u(dst4 + off, tail * sizeof(float), out4);
hvx_vec_store_u(dst5 + off, tail * sizeof(float), out5);
hvx_vec_store_u(dst6 + off, tail * sizeof(float), out6);
hvx_vec_store_u(dst7 + off, tail * sizeof(float), out7);
acc0 = hvx_vec_add_f32_f32(acc0, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out0, vdot), zero));
acc1 = hvx_vec_add_f32_f32(acc1, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out1, vdot), zero));
acc2 = hvx_vec_add_f32_f32(acc2, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out2, vdot), zero));
acc3 = hvx_vec_add_f32_f32(acc3, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out3, vdot), zero));
acc4 = hvx_vec_add_f32_f32(acc4, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out4, vdot), zero));
acc5 = hvx_vec_add_f32_f32(acc5, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out5, vdot), zero));
acc6 = hvx_vec_add_f32_f32(acc6, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out6, vdot), zero));
acc7 = hvx_vec_add_f32_f32(acc7, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out7, vdot), zero));
}
HVX_Vector_x4 accA = { .v = { acc0, acc1, acc2, acc3 } };
HVX_Vector_x4 accB = { .v = { acc4, acc5, acc6, acc7 } };
hvx_vec_store_u(sums + 0, 4 * sizeof(float), hvx_vec_reduce_sum_f32x4(accA));
hvx_vec_store_u(sums + 4, 4 * sizeof(float), hvx_vec_reduce_sum_f32x4(accB));
}
static inline void gdn_mul_scalar_dot8_f32(float * restrict dst0, float * restrict dst1,
float * restrict dst2, float * restrict dst3, float * restrict dst4,
float * restrict dst5, float * restrict dst6, float * restrict dst7,
float mul, const float * restrict dot, uint32_t n, float * restrict sums) {
HVX_Vector acc0 = Q6_V_vzero();
HVX_Vector acc1 = Q6_V_vzero();
HVX_Vector acc2 = Q6_V_vzero();
HVX_Vector acc3 = Q6_V_vzero();
HVX_Vector acc4 = Q6_V_vzero();
HVX_Vector acc5 = Q6_V_vzero();
HVX_Vector acc6 = Q6_V_vzero();
HVX_Vector acc7 = Q6_V_vzero();
const HVX_Vector vmul = hvx_vec_splat_f32(mul);
const uint32_t epv = 128 / sizeof(float);
const uint32_t nvec = n / epv;
const uint32_t tail = n % epv;
for (uint32_t i = 0; i < nvec; ++i) {
HVX_Vector vdot = hvx_vmem(dot + i * epv);
HVX_Vector out0 = hvx_vec_mul_f32_f32(hvx_vmemu(dst0 + i * epv), vmul);
HVX_Vector out1 = hvx_vec_mul_f32_f32(hvx_vmemu(dst1 + i * epv), vmul);
HVX_Vector out2 = hvx_vec_mul_f32_f32(hvx_vmemu(dst2 + i * epv), vmul);
HVX_Vector out3 = hvx_vec_mul_f32_f32(hvx_vmemu(dst3 + i * epv), vmul);
HVX_Vector out4 = hvx_vec_mul_f32_f32(hvx_vmemu(dst4 + i * epv), vmul);
HVX_Vector out5 = hvx_vec_mul_f32_f32(hvx_vmemu(dst5 + i * epv), vmul);
HVX_Vector out6 = hvx_vec_mul_f32_f32(hvx_vmemu(dst6 + i * epv), vmul);
HVX_Vector out7 = hvx_vec_mul_f32_f32(hvx_vmemu(dst7 + i * epv), vmul);
hvx_vmemu(dst0 + i * epv) = out0;
hvx_vmemu(dst1 + i * epv) = out1;
hvx_vmemu(dst2 + i * epv) = out2;
hvx_vmemu(dst3 + i * epv) = out3;
hvx_vmemu(dst4 + i * epv) = out4;
hvx_vmemu(dst5 + i * epv) = out5;
hvx_vmemu(dst6 + i * epv) = out6;
hvx_vmemu(dst7 + i * epv) = out7;
acc0 = hvx_vec_add_f32_f32(acc0, hvx_vec_mul_f32_f32(out0, vdot));
acc1 = hvx_vec_add_f32_f32(acc1, hvx_vec_mul_f32_f32(out1, vdot));
acc2 = hvx_vec_add_f32_f32(acc2, hvx_vec_mul_f32_f32(out2, vdot));
acc3 = hvx_vec_add_f32_f32(acc3, hvx_vec_mul_f32_f32(out3, vdot));
acc4 = hvx_vec_add_f32_f32(acc4, hvx_vec_mul_f32_f32(out4, vdot));
acc5 = hvx_vec_add_f32_f32(acc5, hvx_vec_mul_f32_f32(out5, vdot));
acc6 = hvx_vec_add_f32_f32(acc6, hvx_vec_mul_f32_f32(out6, vdot));
acc7 = hvx_vec_add_f32_f32(acc7, hvx_vec_mul_f32_f32(out7, vdot));
}
if (tail) {
const uint32_t off = nvec * epv;
HVX_Vector vdot = hvx_vmem(dot + off);
HVX_VectorPred mask = Q6_Q_vsetq2_R(tail * sizeof(float));
HVX_Vector zero = Q6_V_vzero();
HVX_Vector out0 = hvx_vec_mul_f32_f32(hvx_vmemu(dst0 + off), vmul);
HVX_Vector out1 = hvx_vec_mul_f32_f32(hvx_vmemu(dst1 + off), vmul);
HVX_Vector out2 = hvx_vec_mul_f32_f32(hvx_vmemu(dst2 + off), vmul);
HVX_Vector out3 = hvx_vec_mul_f32_f32(hvx_vmemu(dst3 + off), vmul);
HVX_Vector out4 = hvx_vec_mul_f32_f32(hvx_vmemu(dst4 + off), vmul);
HVX_Vector out5 = hvx_vec_mul_f32_f32(hvx_vmemu(dst5 + off), vmul);
HVX_Vector out6 = hvx_vec_mul_f32_f32(hvx_vmemu(dst6 + off), vmul);
HVX_Vector out7 = hvx_vec_mul_f32_f32(hvx_vmemu(dst7 + off), vmul);
hvx_vec_store_u(dst0 + off, tail * sizeof(float), out0);
hvx_vec_store_u(dst1 + off, tail * sizeof(float), out1);
hvx_vec_store_u(dst2 + off, tail * sizeof(float), out2);
hvx_vec_store_u(dst3 + off, tail * sizeof(float), out3);
hvx_vec_store_u(dst4 + off, tail * sizeof(float), out4);
hvx_vec_store_u(dst5 + off, tail * sizeof(float), out5);
hvx_vec_store_u(dst6 + off, tail * sizeof(float), out6);
hvx_vec_store_u(dst7 + off, tail * sizeof(float), out7);
acc0 = hvx_vec_add_f32_f32(acc0, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out0, vdot), zero));
acc1 = hvx_vec_add_f32_f32(acc1, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out1, vdot), zero));
acc2 = hvx_vec_add_f32_f32(acc2, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out2, vdot), zero));
acc3 = hvx_vec_add_f32_f32(acc3, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out3, vdot), zero));
acc4 = hvx_vec_add_f32_f32(acc4, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out4, vdot), zero));
acc5 = hvx_vec_add_f32_f32(acc5, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out5, vdot), zero));
acc6 = hvx_vec_add_f32_f32(acc6, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out6, vdot), zero));
acc7 = hvx_vec_add_f32_f32(acc7, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out7, vdot), zero));
}
HVX_Vector_x4 accA = { .v = { acc0, acc1, acc2, acc3 } };
HVX_Vector_x4 accB = { .v = { acc4, acc5, acc6, acc7 } };
hvx_vec_store_u(sums + 0, 4 * sizeof(float), hvx_vec_reduce_sum_f32x4(accA));
hvx_vec_store_u(sums + 4, 4 * sizeof(float), hvx_vec_reduce_sum_f32x4(accB));
}
static inline void gdn_add_scaled_dot8_f32(float * restrict dst0, float * restrict dst1,
float * restrict dst2, float * restrict dst3, float * restrict dst4,
float * restrict dst5, float * restrict dst6, float * restrict dst7,
const float * restrict src, const float * restrict scale,
const float * restrict dot, uint32_t n, float * restrict sums) {
HVX_Vector acc0 = Q6_V_vzero();
HVX_Vector acc1 = Q6_V_vzero();
HVX_Vector acc2 = Q6_V_vzero();
HVX_Vector acc3 = Q6_V_vzero();
HVX_Vector acc4 = Q6_V_vzero();
HVX_Vector acc5 = Q6_V_vzero();
HVX_Vector acc6 = Q6_V_vzero();
HVX_Vector acc7 = Q6_V_vzero();
const HVX_Vector scale0 = hvx_vec_splat_f32(scale[0]);
const HVX_Vector scale1 = hvx_vec_splat_f32(scale[1]);
const HVX_Vector scale2 = hvx_vec_splat_f32(scale[2]);
const HVX_Vector scale3 = hvx_vec_splat_f32(scale[3]);
const HVX_Vector scale4 = hvx_vec_splat_f32(scale[4]);
const HVX_Vector scale5 = hvx_vec_splat_f32(scale[5]);
const HVX_Vector scale6 = hvx_vec_splat_f32(scale[6]);
const HVX_Vector scale7 = hvx_vec_splat_f32(scale[7]);
const uint32_t epv = 128 / sizeof(float);
const uint32_t nvec = n / epv;
const uint32_t tail = n % epv;
for (uint32_t i = 0; i < nvec; ++i) {
HVX_Vector vs = hvx_vmem(src + i * epv);
HVX_Vector vdot = hvx_vmem(dot + i * epv);
HVX_Vector out0 = hvx_vec_add_f32_f32(hvx_vmemu(dst0 + i * epv), hvx_vec_mul_f32_f32(vs, scale0));
HVX_Vector out1 = hvx_vec_add_f32_f32(hvx_vmemu(dst1 + i * epv), hvx_vec_mul_f32_f32(vs, scale1));
HVX_Vector out2 = hvx_vec_add_f32_f32(hvx_vmemu(dst2 + i * epv), hvx_vec_mul_f32_f32(vs, scale2));
HVX_Vector out3 = hvx_vec_add_f32_f32(hvx_vmemu(dst3 + i * epv), hvx_vec_mul_f32_f32(vs, scale3));
HVX_Vector out4 = hvx_vec_add_f32_f32(hvx_vmemu(dst4 + i * epv), hvx_vec_mul_f32_f32(vs, scale4));
HVX_Vector out5 = hvx_vec_add_f32_f32(hvx_vmemu(dst5 + i * epv), hvx_vec_mul_f32_f32(vs, scale5));
HVX_Vector out6 = hvx_vec_add_f32_f32(hvx_vmemu(dst6 + i * epv), hvx_vec_mul_f32_f32(vs, scale6));
HVX_Vector out7 = hvx_vec_add_f32_f32(hvx_vmemu(dst7 + i * epv), hvx_vec_mul_f32_f32(vs, scale7));
hvx_vmemu(dst0 + i * epv) = out0;
hvx_vmemu(dst1 + i * epv) = out1;
hvx_vmemu(dst2 + i * epv) = out2;
hvx_vmemu(dst3 + i * epv) = out3;
hvx_vmemu(dst4 + i * epv) = out4;
hvx_vmemu(dst5 + i * epv) = out5;
hvx_vmemu(dst6 + i * epv) = out6;
hvx_vmemu(dst7 + i * epv) = out7;
acc0 = hvx_vec_add_f32_f32(acc0, hvx_vec_mul_f32_f32(out0, vdot));
acc1 = hvx_vec_add_f32_f32(acc1, hvx_vec_mul_f32_f32(out1, vdot));
acc2 = hvx_vec_add_f32_f32(acc2, hvx_vec_mul_f32_f32(out2, vdot));
acc3 = hvx_vec_add_f32_f32(acc3, hvx_vec_mul_f32_f32(out3, vdot));
acc4 = hvx_vec_add_f32_f32(acc4, hvx_vec_mul_f32_f32(out4, vdot));
acc5 = hvx_vec_add_f32_f32(acc5, hvx_vec_mul_f32_f32(out5, vdot));
acc6 = hvx_vec_add_f32_f32(acc6, hvx_vec_mul_f32_f32(out6, vdot));
acc7 = hvx_vec_add_f32_f32(acc7, hvx_vec_mul_f32_f32(out7, vdot));
}
if (tail) {
const uint32_t off = nvec * epv;
HVX_Vector vs = hvx_vmem(src + off);
HVX_Vector vdot = hvx_vmem(dot + off);
HVX_VectorPred mask = Q6_Q_vsetq2_R(tail * sizeof(float));
HVX_Vector zero = Q6_V_vzero();
HVX_Vector out0 = hvx_vec_add_f32_f32(hvx_vmemu(dst0 + off), hvx_vec_mul_f32_f32(vs, scale0));
HVX_Vector out1 = hvx_vec_add_f32_f32(hvx_vmemu(dst1 + off), hvx_vec_mul_f32_f32(vs, scale1));
HVX_Vector out2 = hvx_vec_add_f32_f32(hvx_vmemu(dst2 + off), hvx_vec_mul_f32_f32(vs, scale2));
HVX_Vector out3 = hvx_vec_add_f32_f32(hvx_vmemu(dst3 + off), hvx_vec_mul_f32_f32(vs, scale3));
HVX_Vector out4 = hvx_vec_add_f32_f32(hvx_vmemu(dst4 + off), hvx_vec_mul_f32_f32(vs, scale4));
HVX_Vector out5 = hvx_vec_add_f32_f32(hvx_vmemu(dst5 + off), hvx_vec_mul_f32_f32(vs, scale5));
HVX_Vector out6 = hvx_vec_add_f32_f32(hvx_vmemu(dst6 + off), hvx_vec_mul_f32_f32(vs, scale6));
HVX_Vector out7 = hvx_vec_add_f32_f32(hvx_vmemu(dst7 + off), hvx_vec_mul_f32_f32(vs, scale7));
hvx_vec_store_u(dst0 + off, tail * sizeof(float), out0);
hvx_vec_store_u(dst1 + off, tail * sizeof(float), out1);
hvx_vec_store_u(dst2 + off, tail * sizeof(float), out2);
hvx_vec_store_u(dst3 + off, tail * sizeof(float), out3);
hvx_vec_store_u(dst4 + off, tail * sizeof(float), out4);
hvx_vec_store_u(dst5 + off, tail * sizeof(float), out5);
hvx_vec_store_u(dst6 + off, tail * sizeof(float), out6);
hvx_vec_store_u(dst7 + off, tail * sizeof(float), out7);
acc0 = hvx_vec_add_f32_f32(acc0, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out0, vdot), zero));
acc1 = hvx_vec_add_f32_f32(acc1, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out1, vdot), zero));
acc2 = hvx_vec_add_f32_f32(acc2, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out2, vdot), zero));
acc3 = hvx_vec_add_f32_f32(acc3, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out3, vdot), zero));
acc4 = hvx_vec_add_f32_f32(acc4, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out4, vdot), zero));
acc5 = hvx_vec_add_f32_f32(acc5, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out5, vdot), zero));
acc6 = hvx_vec_add_f32_f32(acc6, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out6, vdot), zero));
acc7 = hvx_vec_add_f32_f32(acc7, Q6_V_vmux_QVV(mask, hvx_vec_mul_f32_f32(out7, vdot), zero));
}
HVX_Vector_x4 accA = { .v = { acc0, acc1, acc2, acc3 } };
HVX_Vector_x4 accB = { .v = { acc4, acc5, acc6, acc7 } };
hvx_vec_store_u(sums + 0, 4 * sizeof(float), hvx_vec_reduce_sum_f32x4(accA));
hvx_vec_store_u(sums + 4, 4 * sizeof(float), hvx_vec_reduce_sum_f32x4(accB));
}
static void gated_delta_net_f32_pp_thread(unsigned int nth, unsigned int ith, void * data) {
struct htp_gdn_context * gctx = (struct htp_gdn_context *) data;
struct htp_ops_context * octx = gctx->octx;
const struct htp_tensor * q = octx->src[0];
const struct htp_tensor * k = octx->src[1];
const struct htp_tensor * v = octx->src[2];
const struct htp_tensor * g = octx->src[3];
const struct htp_tensor * beta = octx->src[4];
const struct htp_tensor * state = octx->src[5];
const struct htp_tensor * dst = octx->dst;
const uint32_t S_v = v->ne[0];
const uint32_t H = v->ne[1];
const uint32_t n_tokens = v->ne[2];
const uint32_t n_seqs = v->ne[3];
const uint32_t total_rows = H * n_seqs;
if (ith >= total_rows) {
return;
}
const uint32_t rq3 = n_seqs / q->ne[3];
const uint32_t rk3 = n_seqs / k->ne[3];
const float scale = 1.0f / sqrtf((float) S_v);
float * dst_base = (float *) (uintptr_t) dst->data;
float * state_out_base = dst_base + (uint64_t) S_v * H * n_tokens * n_seqs;
const float * state_in_base = (const float *) (uintptr_t) state->data;
const bool kda = (g->ne[0] == S_v);
float local_gate[HTP_GDN_MAX_SV] __attribute__((aligned(128)));
float local_q[HTP_GDN_MAX_SV] __attribute__((aligned(128)));
float local_k[HTP_GDN_MAX_SV] __attribute__((aligned(128)));
float local_sums[4] __attribute__((aligned(128)));
for (uint32_t ir = ith; ir < total_rows; ir += nth) {
const uint32_t iv1 = ir % H;
const uint32_t iv3 = ir / H;
const uint32_t iq1 = iv1 % q->ne[1];
const uint32_t ik1 = iv1 % k->ne[1];
const uint32_t iq3 = iv3 / rq3;
const uint32_t ik3 = iv3 / rk3;
float * s_out = state_out_base + ((uint64_t) iv3 * H + iv1) * S_v * S_v;
const float * s_in = state_in_base + ((uint64_t) iv3 * H + iv1) * S_v * S_v;
memcpy(s_out, s_in, gctx->state_bytes);
float * s_work = s_out;
float * attn_data = dst_base + ((uint64_t) iv3 * n_tokens * H + iv1) * S_v;
for (uint32_t t = 0; t < n_tokens; ++t) {
const float * q_t = (const float *) ((const uint8_t *) (uintptr_t) q->data +
(uint64_t) iq3 * q->nb[3] + (uint64_t) t * q->nb[2] + (uint64_t) iq1 * q->nb[1]);
const float * k_t = (const float *) ((const uint8_t *) (uintptr_t) k->data +
(uint64_t) ik3 * k->nb[3] + (uint64_t) t * k->nb[2] + (uint64_t) ik1 * k->nb[1]);
const float * v_t = (const float *) ((const uint8_t *) (uintptr_t) v->data +
(uint64_t) iv3 * v->nb[3] + (uint64_t) t * v->nb[2] + (uint64_t) iv1 * v->nb[1]);
const float * g_t = (const float *) ((const uint8_t *) (uintptr_t) g->data +
(uint64_t) iv3 * g->nb[3] + (uint64_t) t * g->nb[2] + (uint64_t) iv1 * g->nb[1]);
const float beta_val = *(const float *) ((const uint8_t *) (uintptr_t) beta->data +
(uint64_t) iv3 * beta->nb[3] + (uint64_t) t * beta->nb[2] + (uint64_t) iv1 * beta->nb[1]);
memcpy(local_q, q_t, (size_t) S_v * sizeof(float));
memcpy(local_k, k_t, (size_t) S_v * sizeof(float));
if (kda) {
hvx_exp_f32((uint8_t *) local_gate, (const uint8_t *) g_t, S_v, false);
uint32_t j = 0;
for (; j + 4 <= S_v; j += 4) {
float * row0 = s_work + (uint64_t) (j + 0) * S_v;
float * row1 = s_work + (uint64_t) (j + 1) * S_v;
float * row2 = s_work + (uint64_t) (j + 2) * S_v;
float * row3 = s_work + (uint64_t) (j + 3) * S_v;
gdn_mul_dot4_f32(row0, row1, row2, row3, local_gate, local_k, S_v, local_sums);
float local_delta_b[4] __attribute__((aligned(128)));
for (uint32_t r = 0; r < 4; ++r) {
local_delta_b[r] = (v_t[j + r] - local_sums[r]) * beta_val;
}
gdn_add_scaled_dot4_f32(row0, row1, row2, row3, local_k, local_delta_b, local_q, S_v, local_sums);
for (uint32_t r = 0; r < 4; ++r) {
attn_data[j + r] = local_sums[r] * scale;
}
}
for (; j < S_v; ++j) {
float * row = s_work + (uint64_t) j * S_v;
const float sum = gdn_mul_dot_f32(row, local_gate, local_k, S_v);
const float dj = (v_t[j] - sum) * beta_val;
attn_data[j] = gdn_add_scaled_dot_f32(row, local_k, dj, local_q, S_v) * scale;
}
} else {
const float gate = expf(g_t[0]);
uint32_t j = 0;
for (; j + 4 <= S_v; j += 4) {
float * row0 = s_work + (uint64_t) (j + 0) * S_v;
float * row1 = s_work + (uint64_t) (j + 1) * S_v;
float * row2 = s_work + (uint64_t) (j + 2) * S_v;
float * row3 = s_work + (uint64_t) (j + 3) * S_v;
gdn_mul_scalar_dot4_f32(row0, row1, row2, row3, gate, local_k, S_v, local_sums);
float local_delta_b[4] __attribute__((aligned(128)));
for (uint32_t r = 0; r < 4; ++r) {
local_delta_b[r] = (v_t[j + r] - local_sums[r]) * beta_val;
}
gdn_add_scaled_dot4_f32(row0, row1, row2, row3, local_k, local_delta_b, local_q, S_v, local_sums);
for (uint32_t r = 0; r < 4; ++r) {
attn_data[j + r] = local_sums[r] * scale;
}
}
for (; j < S_v; ++j) {
float * row = s_work + (uint64_t) j * S_v;
const float sum = gdn_mul_scalar_dot_f32(row, gate, local_k, S_v);
const float dj = (v_t[j] - sum) * beta_val;
attn_data[j] = gdn_add_scaled_dot_f32(row, local_k, dj, local_q, S_v) * scale;
}
}
attn_data += (uint64_t) S_v * H;
}
}
}
static void gated_delta_net_f32_tg_thread(unsigned int nth, unsigned int ith, void * data) {
struct htp_gdn_context * gctx = (struct htp_gdn_context *) data;
struct htp_ops_context * octx = gctx->octx;
const struct htp_tensor * q = octx->src[0];
const struct htp_tensor * k = octx->src[1];
const struct htp_tensor * v = octx->src[2];
const struct htp_tensor * g = octx->src[3];
const struct htp_tensor * beta = octx->src[4];
const struct htp_tensor * state = octx->src[5];
const struct htp_tensor * dst = octx->dst;
const uint32_t S_v = v->ne[0];
const uint32_t H = v->ne[1];
const uint32_t n_seqs = v->ne[3];
const uint32_t total_rows = H * n_seqs;
if (ith >= total_rows) {
return;
}
const uint32_t rq3 = n_seqs / q->ne[3];
const uint32_t rk3 = n_seqs / k->ne[3];
const float scale = 1.0f / sqrtf((float) S_v);
float * dst_base = (float *) (uintptr_t) dst->data;
float * state_out_base = dst_base + (uint64_t) S_v * H * n_seqs;
const float * state_in_base = (const float *) (uintptr_t) state->data;
const bool kda = (g->ne[0] == S_v);
float local_gate[HTP_GDN_MAX_SV] __attribute__((aligned(128)));
float local_q[HTP_GDN_MAX_SV] __attribute__((aligned(128)));
float local_k[HTP_GDN_MAX_SV] __attribute__((aligned(128)));
float local_sums[8] __attribute__((aligned(128)));
dma_queue * dma = octx->ctx->dma[ith];
uint8_t * spad = NULL;
if (gctx->use_vtcm) {
spad = gctx->vtcm_state_base + gctx->vtcm_state_per_thread * ith;
}
for (uint32_t ir = ith; ir < total_rows; ir += nth) {
const uint32_t iv1 = ir % H;
const uint32_t iv3 = ir / H;
const uint32_t iq1 = iv1 % q->ne[1];
const uint32_t ik1 = iv1 % k->ne[1];
const uint32_t iq3 = iv3 / rq3;
const uint32_t ik3 = iv3 / rk3;
float * s_out = state_out_base + ((uint64_t) iv3 * H + iv1) * S_v * S_v;
const float * s_in = state_in_base + ((uint64_t) iv3 * H + iv1) * S_v * S_v;
float * s_work;
if (spad) {
dma_queue_push(dma, dma_make_ptr(spad, s_in),
S_v * sizeof(float), S_v * sizeof(float),
S_v * sizeof(float), S_v);
dma_queue_pop(dma);
s_work = (float *) spad;
} else {
s_work = s_out;
memcpy(s_work, s_in, gctx->state_bytes);
}
float * attn_data = dst_base + ((uint64_t) iv3 * H + iv1) * S_v;
const float * q_t = (const float *) ((const uint8_t *) (uintptr_t) q->data +
(uint64_t) iq3 * q->nb[3] + (uint64_t) iq1 * q->nb[1]);
const float * k_t = (const float *) ((const uint8_t *) (uintptr_t) k->data +
(uint64_t) ik3 * k->nb[3] + (uint64_t) ik1 * k->nb[1]);
const float * v_t = (const float *) ((const uint8_t *) (uintptr_t) v->data +
(uint64_t) iv3 * v->nb[3] + (uint64_t) iv1 * v->nb[1]);
const float * g_t = (const float *) ((const uint8_t *) (uintptr_t) g->data +
(uint64_t) iv3 * g->nb[3] + (uint64_t) iv1 * g->nb[1]);
const float beta_val = *(const float *) ((const uint8_t *) (uintptr_t) beta->data +
(uint64_t) iv3 * beta->nb[3] + (uint64_t) iv1 * beta->nb[1]);
memcpy(local_q, q_t, (size_t) S_v * sizeof(float));
memcpy(local_k, k_t, (size_t) S_v * sizeof(float));
if (kda) {
hvx_exp_f32((uint8_t *) local_gate, (const uint8_t *) g_t, S_v, false);
uint32_t j = 0;
for (; j + 8 <= S_v; j += 8) {
float * row0 = s_work + (uint64_t) (j + 0) * S_v;
float * row1 = s_work + (uint64_t) (j + 1) * S_v;
float * row2 = s_work + (uint64_t) (j + 2) * S_v;
float * row3 = s_work + (uint64_t) (j + 3) * S_v;
float * row4 = s_work + (uint64_t) (j + 4) * S_v;
float * row5 = s_work + (uint64_t) (j + 5) * S_v;
float * row6 = s_work + (uint64_t) (j + 6) * S_v;
float * row7 = s_work + (uint64_t) (j + 7) * S_v;
gdn_mul_dot8_f32(row0, row1, row2, row3, row4, row5, row6, row7,
local_gate, local_k, S_v, local_sums);
float local_delta_b[8] __attribute__((aligned(128)));
for (uint32_t r = 0; r < 8; ++r) {
local_delta_b[r] = (v_t[j + r] - local_sums[r]) * beta_val;
}
gdn_add_scaled_dot8_f32(row0, row1, row2, row3, row4, row5, row6, row7,
local_k, local_delta_b, local_q, S_v, local_sums);
for (uint32_t r = 0; r < 8; ++r) {
attn_data[j + r] = local_sums[r] * scale;
}
}
for (; j + 4 <= S_v; j += 4) {
float * row0 = s_work + (uint64_t) (j + 0) * S_v;
float * row1 = s_work + (uint64_t) (j + 1) * S_v;
float * row2 = s_work + (uint64_t) (j + 2) * S_v;
float * row3 = s_work + (uint64_t) (j + 3) * S_v;
gdn_mul_dot4_f32(row0, row1, row2, row3, local_gate, local_k, S_v, local_sums);
float local_delta_b[4] __attribute__((aligned(128)));
for (uint32_t r = 0; r < 4; ++r) {
local_delta_b[r] = (v_t[j + r] - local_sums[r]) * beta_val;
}
gdn_add_scaled_dot4_f32(row0, row1, row2, row3, local_k, local_delta_b, local_q, S_v, local_sums);
for (uint32_t r = 0; r < 4; ++r) {
attn_data[j + r] = local_sums[r] * scale;
}
}
for (; j < S_v; ++j) {
float * row = s_work + (uint64_t) j * S_v;
const float sum = gdn_mul_dot_f32(row, local_gate, local_k, S_v);
const float dj = (v_t[j] - sum) * beta_val;
attn_data[j] = gdn_add_scaled_dot_f32(row, local_k, dj, local_q, S_v) * scale;
}
} else {
const float gate = expf(g_t[0]);
uint32_t j = 0;
for (; j + 8 <= S_v; j += 8) {
float * row0 = s_work + (uint64_t) (j + 0) * S_v;
float * row1 = s_work + (uint64_t) (j + 1) * S_v;
float * row2 = s_work + (uint64_t) (j + 2) * S_v;
float * row3 = s_work + (uint64_t) (j + 3) * S_v;
float * row4 = s_work + (uint64_t) (j + 4) * S_v;
float * row5 = s_work + (uint64_t) (j + 5) * S_v;
float * row6 = s_work + (uint64_t) (j + 6) * S_v;
float * row7 = s_work + (uint64_t) (j + 7) * S_v;
gdn_mul_scalar_dot8_f32(row0, row1, row2, row3, row4, row5, row6, row7,
gate, local_k, S_v, local_sums);
float local_delta_b[8] __attribute__((aligned(128)));
for (uint32_t r = 0; r < 8; ++r) {
local_delta_b[r] = (v_t[j + r] - local_sums[r]) * beta_val;
}
gdn_add_scaled_dot8_f32(row0, row1, row2, row3, row4, row5, row6, row7,
local_k, local_delta_b, local_q, S_v, local_sums);
for (uint32_t r = 0; r < 8; ++r) {
attn_data[j + r] = local_sums[r] * scale;
}
}
for (; j + 4 <= S_v; j += 4) {
float * row0 = s_work + (uint64_t) (j + 0) * S_v;
float * row1 = s_work + (uint64_t) (j + 1) * S_v;
float * row2 = s_work + (uint64_t) (j + 2) * S_v;
float * row3 = s_work + (uint64_t) (j + 3) * S_v;
gdn_mul_scalar_dot4_f32(row0, row1, row2, row3, gate, local_k, S_v, local_sums);
float local_delta_b[4] __attribute__((aligned(128)));
for (uint32_t r = 0; r < 4; ++r) {
local_delta_b[r] = (v_t[j + r] - local_sums[r]) * beta_val;
}
gdn_add_scaled_dot4_f32(row0, row1, row2, row3, local_k, local_delta_b, local_q, S_v, local_sums);
for (uint32_t r = 0; r < 4; ++r) {
attn_data[j + r] = local_sums[r] * scale;
}
}
for (; j < S_v; ++j) {
float * row = s_work + (uint64_t) j * S_v;
const float sum = gdn_mul_scalar_dot_f32(row, gate, local_k, S_v);
const float dj = (v_t[j] - sum) * beta_val;
attn_data[j] = gdn_add_scaled_dot_f32(row, local_k, dj, local_q, S_v) * scale;
}
}
if (spad) {
dma_queue_push(dma, dma_make_ptr(s_out, spad),
S_v * sizeof(float), S_v * sizeof(float),
S_v * sizeof(float), S_v);
dma_queue_pop(dma);
}
}
}
int op_gated_delta_net(struct htp_ops_context * octx) {
const struct htp_tensor * q = octx->src[0];
const struct htp_tensor * k = octx->src[1];
const struct htp_tensor * v = octx->src[2];
const struct htp_tensor * g = octx->src[3];
const struct htp_tensor * beta = octx->src[4];
const struct htp_tensor * state = octx->src[5];
const struct htp_tensor * dst = octx->dst;
if (!q || !k || !v || !g || !beta || !state || !dst) {
return HTP_STATUS_INVAL_PARAMS;
}
if (q->type != HTP_TYPE_F32 || k->type != HTP_TYPE_F32 || v->type != HTP_TYPE_F32 ||
g->type != HTP_TYPE_F32 || beta->type != HTP_TYPE_F32 || state->type != HTP_TYPE_F32 ||
dst->type != HTP_TYPE_F32) {
return HTP_STATUS_NO_SUPPORT;
}
const uint32_t S_v = v->ne[0];
const uint32_t H = v->ne[1];
const uint32_t n_tokens = v->ne[2];
const uint32_t n_seqs = v->ne[3];
if (S_v == 0 || S_v > HTP_GDN_MAX_SV || H == 0 || n_tokens == 0 || n_seqs == 0) {
return HTP_STATUS_NO_SUPPORT;
}
if ((g->ne[0] != 1 && g->ne[0] != S_v) || beta->ne[0] != 1) {
return HTP_STATUS_NO_SUPPORT;
}
if (q->ne[0] != S_v || k->ne[0] != S_v || q->ne[1] == 0 || k->ne[1] == 0 ||
q->ne[2] != n_tokens || k->ne[2] != n_tokens || q->ne[3] == 0 || k->ne[3] == 0 ||
(n_seqs % q->ne[3]) != 0 || (n_seqs % k->ne[3]) != 0) {
return HTP_STATUS_NO_SUPPORT;
}
if (state->ne[0] * state->ne[1] * state->ne[2] * state->ne[3] != S_v * S_v * H * n_seqs) {
return HTP_STATUS_NO_SUPPORT;
}
if (dst->ne[0] != S_v * H || dst->ne[1] != n_tokens * n_seqs + S_v * n_seqs) {
return HTP_STATUS_NO_SUPPORT;
}
if (octx->flags & HTP_OPFLAGS_SKIP_COMPUTE) {
return HTP_STATUS_OK;
}
struct htp_gdn_context gctx;
gctx.octx = octx;
gctx.rows_per_thread = (H * n_seqs + octx->n_threads - 1) / octx->n_threads;
gctx.state_bytes = (size_t) S_v * S_v * sizeof(float);
size_t state_aligned = (size_t) S_v * S_v * sizeof(float);
state_aligned = (state_aligned + 127) & ~(size_t)127;
gctx.use_vtcm = false;
gctx.vtcm_state_base = NULL;
gctx.vtcm_state_per_thread = 0;
if (n_tokens == 1 && octx->ctx->vtcm_base) {
size_t vtcm_total = state_aligned * octx->n_threads;
if (octx->ctx->vtcm_size >= vtcm_total) {
gctx.use_vtcm = true;
gctx.vtcm_state_base = octx->ctx->vtcm_base;
gctx.vtcm_state_per_thread = state_aligned;
}
}
if (n_tokens == 1) {
worker_pool_run_func(octx->ctx->worker_pool, gated_delta_net_f32_tg_thread, &gctx, octx->n_threads);
} else {
worker_pool_run_func(octx->ctx->worker_pool, gated_delta_net_f32_pp_thread, &gctx, octx->n_threads);
}
return HTP_STATUS_OK;
}

View file

@ -87,17 +87,17 @@ static void ggml_backend_metal_buffer_shared_clear(ggml_backend_buffer_t buffer,
}
static ggml_backend_buffer_i ggml_backend_metal_buffer_shared_i = {
/* .free_buffer = */ ggml_backend_metal_buffer_shared_free_buffer,
/* .get_base = */ ggml_backend_metal_buffer_shared_get_base,
/* .init_tensor = */ NULL,
/* .memset_tensor = */ ggml_backend_metal_buffer_shared_memset_tensor,
/* .set_tensor = */ ggml_backend_metal_buffer_shared_set_tensor,
/* .get_tensor = */ ggml_backend_metal_buffer_shared_get_tensor,
/* .set_tensor_2d = */ NULL,
/* .get_tensor_2d = */ NULL,
/* .cpy_tensor = */ ggml_backend_metal_buffer_shared_cpy_tensor,
/* .clear = */ ggml_backend_metal_buffer_shared_clear,
/* .reset = */ NULL,
/* .free_buffer = */ ggml_backend_metal_buffer_shared_free_buffer,
/* .get_base = */ ggml_backend_metal_buffer_shared_get_base,
/* .init_tensor = */ NULL,
/* .memset_tensor = */ ggml_backend_metal_buffer_shared_memset_tensor,
/* .set_tensor = */ ggml_backend_metal_buffer_shared_set_tensor,
/* .get_tensor = */ ggml_backend_metal_buffer_shared_get_tensor,
/* .set_tensor_2d = */ NULL,
/* .get_tensor_2d = */ NULL,
/* .cpy_tensor = */ ggml_backend_metal_buffer_shared_cpy_tensor,
/* .clear = */ ggml_backend_metal_buffer_shared_clear,
/* .reset = */ NULL,
};
// private buffer
@ -163,17 +163,17 @@ static void ggml_backend_metal_buffer_private_clear(ggml_backend_buffer_t buffer
}
static ggml_backend_buffer_i ggml_backend_metal_buffer_private_i = {
/* .free_buffer = */ ggml_backend_metal_buffer_private_free_buffer,
/* .get_base = */ ggml_backend_metal_buffer_private_get_base,
/* .init_tensor = */ NULL,
/* .memset_tensor = */ ggml_backend_metal_buffer_private_memset_tensor,
/* .set_tensor = */ ggml_backend_metal_buffer_private_set_tensor,
/* .get_tensor = */ ggml_backend_metal_buffer_private_get_tensor,
/* .set_tensor_2d = */ NULL,
/* .get_tensor_2d = */ NULL,
/* .cpy_tensor = */ ggml_backend_metal_buffer_private_cpy_tensor,
/* .clear = */ ggml_backend_metal_buffer_private_clear,
/* .reset = */ NULL,
/* .free_buffer = */ ggml_backend_metal_buffer_private_free_buffer,
/* .get_base = */ ggml_backend_metal_buffer_private_get_base,
/* .init_tensor = */ NULL,
/* .memset_tensor = */ ggml_backend_metal_buffer_private_memset_tensor,
/* .set_tensor = */ ggml_backend_metal_buffer_private_set_tensor,
/* .get_tensor = */ ggml_backend_metal_buffer_private_get_tensor,
/* .set_tensor_2d = */ NULL,
/* .get_tensor_2d = */ NULL,
/* .cpy_tensor = */ ggml_backend_metal_buffer_private_cpy_tensor,
/* .clear = */ ggml_backend_metal_buffer_private_clear,
/* .reset = */ NULL,
};
static bool ggml_backend_buffer_is_metal(ggml_backend_buffer_t buffer) {

View file

@ -0,0 +1,252 @@
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
#pragma OPENCL EXTENSION cl_qcom_subgroup_uniform_load: enable
#pragma OPENCL EXTENSION cl_qcom_subgroup_constant_load: enable
#pragma OPENCL EXTENSION cl_qcom_extra_vector_types : enable
#define TILESIZE_K 16
#define TILESIZE_M 64
#define TILESIZE_N 32
#define dequantize_q4_0(q4, a_f16, scale) \
a_f16.s0 = (half)((q4.s0 & 0x000F) - 8) * scale; \
a_f16.s1 = (half)(((q4.s0 & 0x00F0) >> 4) - 8) * scale; \
a_f16.s2 = (half)(((q4.s0 & 0x0F00) >> 8) - 8) * scale; \
a_f16.s3 = (half)(((q4.s0 & 0xF000) >> 12) - 8) * scale; \
a_f16.s4 = (half)((q4.s1 & 0x000F) - 8) * scale; \
a_f16.s5 = (half)(((q4.s1 & 0x00F0) >> 4) - 8) * scale; \
a_f16.s6 = (half)(((q4.s1 & 0x0F00) >> 8) - 8) * scale; \
a_f16.s7 = (half)(((q4.s1 & 0xF000) >> 12) - 8) * scale; \
a_f16.s8 = (half)((q4.s2 & 0x000F) - 8) * scale; \
a_f16.s9 = (half)(((q4.s2 & 0x00F0) >> 4) - 8) * scale; \
a_f16.sa = (half)(((q4.s2 & 0x0F00) >> 8) - 8) * scale; \
a_f16.sb = (half)(((q4.s2 & 0xF000) >> 12) - 8) * scale; \
a_f16.sc = (half)((q4.s3 & 0x000F) - 8) * scale; \
a_f16.sd = (half)(((q4.s3 & 0x00F0) >> 4) - 8) * scale; \
a_f16.se = (half)(((q4.s3 & 0x0F00) >> 8) - 8) * scale; \
a_f16.sf = (half)(((q4.s3 & 0xF000) >> 12) - 8) * scale; \
#define dotx16_reduce8(a_reg, b_lm, c_reg, lm_offset) \
acc.s0 = dot(a_reg.s0123, b_lm[lm_offset + 0]); \
acc.s1 = dot(a_reg.s0123, b_lm[lm_offset + 1]); \
acc.s2 = dot(a_reg.s0123, b_lm[lm_offset + 2]); \
acc.s3 = dot(a_reg.s0123, b_lm[lm_offset + 3]); \
acc.s4 = dot(a_reg.s0123, b_lm[lm_offset + 4]); \
acc.s5 = dot(a_reg.s0123, b_lm[lm_offset + 5]); \
acc.s6 = dot(a_reg.s0123, b_lm[lm_offset + 6]); \
acc.s7 = dot(a_reg.s0123, b_lm[lm_offset + 7]); \
acc.s8 = dot(a_reg.s0123, b_lm[lm_offset + 8]); \
acc.s9 = dot(a_reg.s0123, b_lm[lm_offset + 9]); \
acc.sa = dot(a_reg.s0123, b_lm[lm_offset + 10]); \
acc.sb = dot(a_reg.s0123, b_lm[lm_offset + 11]); \
acc.sc = dot(a_reg.s0123, b_lm[lm_offset + 12]); \
acc.sd = dot(a_reg.s0123, b_lm[lm_offset + 13]); \
acc.se = dot(a_reg.s0123, b_lm[lm_offset + 14]); \
acc.sf = dot(a_reg.s0123, b_lm[lm_offset + 15]); \
acc.s0 += dot(a_reg.s4567, b_lm[lm_offset + 32]); \
acc.s1 += dot(a_reg.s4567, b_lm[lm_offset + 33]); \
acc.s2 += dot(a_reg.s4567, b_lm[lm_offset + 34]); \
acc.s3 += dot(a_reg.s4567, b_lm[lm_offset + 35]); \
acc.s4 += dot(a_reg.s4567, b_lm[lm_offset + 36]); \
acc.s5 += dot(a_reg.s4567, b_lm[lm_offset + 37]); \
acc.s6 += dot(a_reg.s4567, b_lm[lm_offset + 38]); \
acc.s7 += dot(a_reg.s4567, b_lm[lm_offset + 39]); \
acc.s8 += dot(a_reg.s4567, b_lm[lm_offset + 40]); \
acc.s9 += dot(a_reg.s4567, b_lm[lm_offset + 41]); \
acc.sa += dot(a_reg.s4567, b_lm[lm_offset + 42]); \
acc.sb += dot(a_reg.s4567, b_lm[lm_offset + 43]); \
acc.sc += dot(a_reg.s4567, b_lm[lm_offset + 44]); \
acc.sd += dot(a_reg.s4567, b_lm[lm_offset + 45]); \
acc.se += dot(a_reg.s4567, b_lm[lm_offset + 46]); \
acc.sf += dot(a_reg.s4567, b_lm[lm_offset + 47]); \
c_reg.lo += convert_float8(acc.lo); \
c_reg.hi += convert_float8(acc.hi); \
acc.s0 = dot(a_reg.s89ab, b_lm[lm_offset + 64]); \
acc.s1 = dot(a_reg.s89ab, b_lm[lm_offset + 65]); \
acc.s2 = dot(a_reg.s89ab, b_lm[lm_offset + 66]); \
acc.s3 = dot(a_reg.s89ab, b_lm[lm_offset + 67]); \
acc.s4 = dot(a_reg.s89ab, b_lm[lm_offset + 68]); \
acc.s5 = dot(a_reg.s89ab, b_lm[lm_offset + 69]); \
acc.s6 = dot(a_reg.s89ab, b_lm[lm_offset + 70]); \
acc.s7 = dot(a_reg.s89ab, b_lm[lm_offset + 71]); \
acc.s8 = dot(a_reg.s89ab, b_lm[lm_offset + 72]); \
acc.s9 = dot(a_reg.s89ab, b_lm[lm_offset + 73]); \
acc.sa = dot(a_reg.s89ab, b_lm[lm_offset + 74]); \
acc.sb = dot(a_reg.s89ab, b_lm[lm_offset + 75]); \
acc.sc = dot(a_reg.s89ab, b_lm[lm_offset + 76]); \
acc.sd = dot(a_reg.s89ab, b_lm[lm_offset + 77]); \
acc.se = dot(a_reg.s89ab, b_lm[lm_offset + 78]); \
acc.sf = dot(a_reg.s89ab, b_lm[lm_offset + 79]); \
acc.s0 += dot(a_reg.scdef, b_lm[lm_offset + 96]); \
acc.s1 += dot(a_reg.scdef, b_lm[lm_offset + 97]); \
acc.s2 += dot(a_reg.scdef, b_lm[lm_offset + 98]); \
acc.s3 += dot(a_reg.scdef, b_lm[lm_offset + 99]); \
acc.s4 += dot(a_reg.scdef, b_lm[lm_offset + 100]); \
acc.s5 += dot(a_reg.scdef, b_lm[lm_offset + 101]); \
acc.s6 += dot(a_reg.scdef, b_lm[lm_offset + 102]); \
acc.s7 += dot(a_reg.scdef, b_lm[lm_offset + 103]); \
acc.s8 += dot(a_reg.scdef, b_lm[lm_offset + 104]); \
acc.s9 += dot(a_reg.scdef, b_lm[lm_offset + 105]); \
acc.sa += dot(a_reg.scdef, b_lm[lm_offset + 106]); \
acc.sb += dot(a_reg.scdef, b_lm[lm_offset + 107]); \
acc.sc += dot(a_reg.scdef, b_lm[lm_offset + 108]); \
acc.sd += dot(a_reg.scdef, b_lm[lm_offset + 109]); \
acc.se += dot(a_reg.scdef, b_lm[lm_offset + 110]); \
acc.sf += dot(a_reg.scdef, b_lm[lm_offset + 111]); \
c_reg.lo += convert_float8(acc.lo); \
c_reg.hi += convert_float8(acc.hi); \
__attribute__((qcom_wave_pair_mode(1))) // 1=force single 2=force pair
kernel void kernel_gemm_moe_q4_0_f32_ns(
__read_only image1d_buffer_t src0_q,
__global half * src0_d,
__read_only image1d_buffer_t src1,
__global uint * src2,
__global ushort * src2_emap,
__write_only image1d_buffer_t dst,
__global int * total_tiles,
uint ne00,
uint ne01
) {
uint block_id_m = get_global_id(1); // m_tile
uint block_id_n = get_global_id(2); // n_tile
// Boundary check
if (((get_global_id(0) + block_id_m * TILESIZE_M) >= ne01) || (block_id_n >= total_tiles[0])) {
return;
}
__private half16 reg_a;
__private float32 reg_c = (float32)(0);
__local half4 shared_b[128];
const ushort expert_id = src2_emap[block_id_n];
const uint row = block_id_m * TILESIZE_M;
const uint col = block_id_n * TILESIZE_N;
uint sub_block_id_m = get_local_id(0);
uint2 b_global_offset;
b_global_offset.x = ((sub_block_id_m & 3) << 2) + (sub_block_id_m >> 2) * ne00;
b_global_offset.y = b_global_offset.x + (16 * ne00);
uint2 b_local_offset;
b_local_offset.x = (sub_block_id_m & 3) * 32 + (sub_block_id_m >> 2);
b_local_offset.y = b_local_offset.x + 16;
// Loop along K axis, 32 elements (one block) for each iteration, divided into 2 sub-blocks
for (uint step = 0; step < ne00; step += TILESIZE_K * 2) {
// First sub-block
uint q_sub_offset = row + ((ne01 * step) >> 3) + ((expert_id * ne00 * ne01) >> 3);
uint s_sub_offset = row + ((ne01 * step) >> 5) + ((expert_id * ne00 * ne01) >> 5);
uint b_sub_offset = col * ne00 + step;
// Load scale for current Q4_0 block
uint s_offset = s_sub_offset + get_global_id(0);
half s = src0_d[s_offset];
// Load 16 q (64-bits) in transposed layout
uint2 q4x16;
q4x16.x = read_imageui(src0_q, q_sub_offset + sub_block_id_m).x;
q4x16.y = read_imageui(src0_q, q_sub_offset + sub_block_id_m + ne01).x;
// Load 16x32 floats from matrix B, each fiber out of 64 in a sub-group loads 8 elements
float8 bx8_f32;
bx8_f32.lo = read_imagef(src1, (b_sub_offset + b_global_offset.x) / 4);
bx8_f32.hi = read_imagef(src1, (b_sub_offset + b_global_offset.y) / 4);
// Convert to half and store to LM to share within the subgroup
half8 bx8_f16 = convert_half8(bx8_f32);
shared_b[b_local_offset.x] = bx8_f16.lo;
shared_b[b_local_offset.y] = bx8_f16.hi;
// Dequantization
dequantize_q4_0(as_ushort4(q4x16), reg_a, s);
sub_group_barrier(CLK_LOCAL_MEM_FENCE);
// 32 16x16 fp16 dot product with 8 elements reduction for better precision
half16 acc;
dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
// Repeat for second sub-block
uint half_step = step + TILESIZE_K;
q_sub_offset = row + ((ne01 * half_step) >> 3) + ((expert_id * ne00 * ne01) >> 3);
b_sub_offset = col * ne00 + half_step;
// Load next 16 q (64-bits) in transposed layout
q4x16.x = read_imageui(src0_q, q_sub_offset + sub_block_id_m).x;
q4x16.y = read_imageui(src0_q, q_sub_offset + sub_block_id_m + ne01).x;
// Load 16x32 floats from matrix B, each fiber out of 64 in a sub-group loads 8 elements
bx8_f32.lo = read_imagef(src1, (b_sub_offset + b_global_offset.x) / 4);
bx8_f32.hi = read_imagef(src1, (b_sub_offset + b_global_offset.y) / 4);
// Convert to half and store to LM to share within the subgroup
bx8_f16 = convert_half8(bx8_f32);
shared_b[b_local_offset.x] = bx8_f16.lo;
shared_b[b_local_offset.y] = bx8_f16.hi;
// Dequantization
dequantize_q4_0(as_ushort4(q4x16), reg_a, s);
sub_group_barrier(CLK_LOCAL_MEM_FENCE);
// 32 16x16 fp16 dot product with 3-levels reduction for better precision
dotx16_reduce8(reg_a, shared_b, reg_c.lo, 0);
dotx16_reduce8(reg_a, shared_b, reg_c.hi, 16);
}
// Load poster router and share in LM
__local uint out_idx[TILESIZE_N];
if (get_local_id(0) < TILESIZE_N) {
uint idx = src2[block_id_n * TILESIZE_N + get_local_id(0)];
if (idx == 0xFFFFFFFF) {
idx = src2[block_id_n * TILESIZE_N + 0];
}
out_idx[get_local_id(0)] = idx * ne01;
}
barrier(CLK_LOCAL_MEM_FENCE);
// Scatter results back to original position in output grid
uint m_offset = row + get_local_id(0);
write_imagef(dst, out_idx[1] + m_offset, (reg_c.s1));
write_imagef(dst, out_idx[2] + m_offset, (reg_c.s2));
write_imagef(dst, out_idx[3] + m_offset, (reg_c.s3));
write_imagef(dst, out_idx[4] + m_offset, (reg_c.s4));
write_imagef(dst, out_idx[5] + m_offset, (reg_c.s5));
write_imagef(dst, out_idx[6] + m_offset, (reg_c.s6));
write_imagef(dst, out_idx[7] + m_offset, (reg_c.s7));
write_imagef(dst, out_idx[8] + m_offset, (reg_c.s8));
write_imagef(dst, out_idx[9] + m_offset, (reg_c.s9));
write_imagef(dst, out_idx[10] + m_offset, (reg_c.sa));
write_imagef(dst, out_idx[11] + m_offset, (reg_c.sb));
write_imagef(dst, out_idx[12] + m_offset, (reg_c.sc));
write_imagef(dst, out_idx[13] + m_offset, (reg_c.sd));
write_imagef(dst, out_idx[14] + m_offset, (reg_c.se));
write_imagef(dst, out_idx[15] + m_offset, (reg_c.sf));
write_imagef(dst, out_idx[16] + m_offset, (reg_c.sg));
write_imagef(dst, out_idx[17] + m_offset, (reg_c.sh));
write_imagef(dst, out_idx[18] + m_offset, (reg_c.si));
write_imagef(dst, out_idx[19] + m_offset, (reg_c.sj));
write_imagef(dst, out_idx[20] + m_offset, (reg_c.sk));
write_imagef(dst, out_idx[21] + m_offset, (reg_c.sl));
write_imagef(dst, out_idx[22] + m_offset, (reg_c.sm));
write_imagef(dst, out_idx[23] + m_offset, (reg_c.sn));
write_imagef(dst, out_idx[24] + m_offset, (reg_c.so));
write_imagef(dst, out_idx[25] + m_offset, (reg_c.sp));
write_imagef(dst, out_idx[26] + m_offset, (reg_c.sq));
write_imagef(dst, out_idx[27] + m_offset, (reg_c.sr));
write_imagef(dst, out_idx[28] + m_offset, (reg_c.ss));
write_imagef(dst, out_idx[29] + m_offset, (reg_c.st));
write_imagef(dst, out_idx[30] + m_offset, (reg_c.su));
write_imagef(dst, out_idx[31] + m_offset, (reg_c.sv));
// Store zero padding parts to the index of first output in tile, override correct result in the end
barrier(CLK_GLOBAL_MEM_FENCE);
write_imagef(dst, out_idx[0] + m_offset, (reg_c.s0));
}

View file

@ -0,0 +1,116 @@
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
#define QK_Q4_0 32
#define N_SIMDGROUP 4
#define SIMDGROUP_WIDTH 64
static inline float8 q4_0_to_fp32_packed8(ushort2 q4x8) {
float8 fp32x8;
fp32x8.s0 = (float)((q4x8.s0 & 0x000F) - 8);
fp32x8.s1 = (float)(((q4x8.s0 & 0x00F0) >> 4) - 8);
fp32x8.s2 = (float)(((q4x8.s0 & 0x0F00) >> 8) - 8);
fp32x8.s3 = (float)(((q4x8.s0 & 0xF000) >> 12) - 8);
fp32x8.s4 = (float)((q4x8.s1 & 0x000F) - 8);
fp32x8.s5 = (float)(((q4x8.s1 & 0x00F0) >> 4) - 8);
fp32x8.s6 = (float)(((q4x8.s1 & 0x0F00) >> 8) - 8);
fp32x8.s7 = (float)(((q4x8.s1 & 0xF000) >> 12) - 8);
return fp32x8;
}
__attribute__((qcom_reqd_sub_group_size("half")))
__kernel void kernel_gemv_moe_q4_0_f32_ns(
__global uint * src0_q,
__global half * src0_d,
__read_only image1d_buffer_t src1,
__global uint * src2,
__global float * dst,
ulong offsetd,
int ne00,
int ne01,
int ne11
) {
uint i01 = get_global_id(0);
uint i20 = get_global_id(2);
uint sgid = get_local_id(1);
uint slid = get_sub_group_local_id();
uint i11 = i20 % ne11;
uint expert_id = src2[i20];
uint expert_offset = expert_id * ne00 * ne01 / 32;
__private float sum = 0.0f; // each thread calculate partial sum of one output
// loop along ne00 in block granularity, skip 4 blocks every iter
for (uint ib00 = sgid; ib00 < (ne00 / QK_Q4_0); ib00 += N_SIMDGROUP) {
// load one block of q
uint4 regQ;
uint block_offset = expert_offset * 4 + ib00 * ne01 * 4 + i01;
regQ.s0 = src0_q[block_offset];
regQ.s1 = src0_q[block_offset + ne01];
regQ.s2 = src0_q[block_offset + ne01 * 2];
regQ.s3 = src0_q[block_offset + ne01 * 3];
uint offset = i11 * ne00 / 4 + ib00 * 8;
float8 fp32x8 = q4_0_to_fp32_packed8(as_ushort2(regQ.s0));
float4 shared_y4;
shared_y4 = read_imagef(src1, (offset + 0));
float4 acc = shared_y4 * fp32x8.lo;
shared_y4 = read_imagef(src1, (offset + 1));
acc += shared_y4 * fp32x8.hi;
fp32x8 = q4_0_to_fp32_packed8(as_ushort2(regQ.s1));
shared_y4 = read_imagef(src1, (offset + 2));
acc += shared_y4 * fp32x8.lo;
shared_y4 = read_imagef(src1, (offset + 3));
acc += shared_y4 * fp32x8.hi;
fp32x8 = q4_0_to_fp32_packed8(as_ushort2(regQ.s2));
shared_y4 = read_imagef(src1, (offset + 4));
acc += shared_y4 * fp32x8.lo;
shared_y4 = read_imagef(src1, (offset + 5));
acc += shared_y4 * fp32x8.hi;
fp32x8 = q4_0_to_fp32_packed8(as_ushort2(regQ.s3));
shared_y4 = read_imagef(src1, (offset + 6));
acc += shared_y4 * fp32x8.lo;
shared_y4 = read_imagef(src1, (offset + 7));
acc += shared_y4 * fp32x8.hi;
half regS = src0_d[ib00 * ne01 + i01 + expert_offset];
sum += (float)(regS) * ((acc.s0 + acc.s1) + (acc.s2 + acc.s3));
}
// reduction in local memory, assumes #subgroups=4
__local float reduceLM[SIMDGROUP_WIDTH * (N_SIMDGROUP - 1)];
if (sgid == 1) reduceLM[SIMDGROUP_WIDTH * 0 + slid] = sum;
if (sgid == 2) reduceLM[SIMDGROUP_WIDTH * 1 + slid] = sum;
if (sgid == 3) reduceLM[SIMDGROUP_WIDTH * 2 + slid] = sum;
barrier(CLK_LOCAL_MEM_FENCE);
if (sgid == 0) sum += reduceLM[SIMDGROUP_WIDTH * 0 + slid];
if (sgid == 0) sum += reduceLM[SIMDGROUP_WIDTH * 1 + slid];
if (sgid == 0) sum += reduceLM[SIMDGROUP_WIDTH * 2 + slid];
// 1 outputs per thread in subgroup 0
if (sgid == 0) {
dst = dst + (offsetd >> 2);
dst[i01 + i20 * ne01] = sum;
}
}

View file

@ -0,0 +1,148 @@
#include "cumsum.hpp"
#include "common.hpp"
#include <algorithm>
#define SYCL_CUMSUM_BLOCK_SIZE 256
static __dpct_inline__ float warp_prefix_inclusive_sum_f32(float x, const sycl::nd_item<3> & item) {
return sycl::inclusive_scan_over_group(item.get_sub_group(), x, sycl::plus<float>());
}
static void cumsum_f32_kernel(
const float * __restrict__ src, float * __restrict__ dst,
const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03,
const int64_t s01, const int64_t s02, const int64_t s03,
const int64_t d1, const int64_t d2, const int64_t d3,
const sycl::nd_item<3> & item, float * smem) {
const int tid = item.get_local_id(2);
const int block_size = item.get_local_range(2);
const int lane = tid % WARP_SIZE;
const int warp = tid / WARP_SIZE;
const int warps_per_block = block_size / WARP_SIZE;
float * s_vals = smem;
float * s_warp_sums = smem + block_size;
float * s_carry = smem + block_size + warps_per_block;
if (tid == 0) {
s_carry[0] = 0.0f;
}
item.barrier(sycl::access::fence_space::local_space);
const int64_t i3 = item.get_group(0);
const int64_t i2 = item.get_group(1);
const int64_t i1 = item.get_group(2);
if (i3 >= ne03 || i2 >= ne02 || i1 >= ne01) {
return;
}
const float * src_row = src + i1 * s01 + i2 * s02 + i3 * s03;
float * dst_row = dst + i1 * d1 + i2 * d2 + i3 * d3;
constexpr int num_unroll = 4;
float temp[num_unroll];
for (int64_t i = 0; i < ne00; i += num_unroll * block_size) {
int64_t idx = i + tid * num_unroll;
temp[0] = (idx < ne00 ? src_row[idx] : 0.0f);
#pragma unroll
for (int j = 1; j < num_unroll; j++) {
temp[j] = temp[j - 1];
if (idx + j < ne00) {
temp[j] += src_row[idx + j];
}
}
float val = (idx < ne00) ? temp[num_unroll - 1] : 0.0f;
val = warp_prefix_inclusive_sum_f32(val, item);
s_vals[tid] = val;
if (lane == WARP_SIZE - 1) {
s_warp_sums[warp] = val;
}
item.barrier(sycl::access::fence_space::local_space);
if (warp == 0) {
float w = (tid < warps_per_block) ? s_warp_sums[tid] : 0.0f;
float inc = warp_prefix_inclusive_sum_f32(w, item);
if (tid < warps_per_block) {
s_warp_sums[tid] = inc - w;
}
if (tid == warps_per_block - 1) {
s_carry[1] = inc;
}
}
item.barrier(sycl::access::fence_space::local_space);
float carry = s_carry[0];
float final_offset = s_vals[tid] + s_warp_sums[warp] + carry - temp[num_unroll - 1];
#pragma unroll
for (int j = 0; j < num_unroll; j++) {
if (idx + j < ne00) {
dst_row[idx + j] = temp[j] + final_offset;
}
}
item.barrier(sycl::access::fence_space::local_space);
if (tid == 0) {
s_carry[0] += s_carry[1];
}
}
}
inline void ggml_sycl_op_cumsum(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
const ggml_tensor * src0 = dst->src[0];
GGML_ASSERT(src0->type == GGML_TYPE_F32);
GGML_ASSERT(dst->type == GGML_TYPE_F32);
dpct::queue_ptr stream = ctx.stream();
SYCL_CHECK(ggml_sycl_set_device(ctx.device));
const float * src_d = static_cast<const float *>(src0->data);
float * dst_d = static_cast<float *>(dst->data);
const int64_t ne00 = src0->ne[0];
const int64_t ne01 = src0->ne[1];
const int64_t ne02 = src0->ne[2];
const int64_t ne03 = src0->ne[3];
const size_t ts = sizeof(float);
const int64_t s01 = src0->nb[1] / ts;
const int64_t s02 = src0->nb[2] / ts;
const int64_t s03 = src0->nb[3] / ts;
const int64_t d1 = dst->nb[1] / ts;
const int64_t d2 = dst->nb[2] / ts;
const int64_t d3 = dst->nb[3] / ts;
const int num_warps = (ne00 + WARP_SIZE - 1) / WARP_SIZE;
int block_size = num_warps * WARP_SIZE;
block_size = std::min(block_size, SYCL_CUMSUM_BLOCK_SIZE);
const int warps_per_block = block_size / WARP_SIZE;
const int smem_size = block_size + warps_per_block + 2;
const sycl::range<3> grid(ne03, ne02, ne01);
const sycl::range<3> block(1, 1, block_size);
stream->submit([&](sycl::handler & cgh) {
sycl::local_accessor<float, 1> smem_acc(sycl::range<1>(smem_size), cgh);
cgh.parallel_for(
sycl::nd_range<3>(grid * block, block),
[=](sycl::nd_item<3> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
cumsum_f32_kernel(src_d, dst_d, ne00, ne01, ne02, ne03,
s01, s02, s03, d1, d2, d3,
item, get_pointer(smem_acc));
});
});
}
void ggml_sycl_cumsum(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
ggml_sycl_op_cumsum(ctx, dst);
}

View file

@ -0,0 +1,5 @@
#pragma once
#include "common.hpp"
void ggml_sycl_cumsum(ggml_backend_sycl_context & ctx, ggml_tensor * dst);

View file

@ -0,0 +1,67 @@
#include "diag.hpp"
#include "common.hpp"
#define SYCL_DIAG_BLOCK_SIZE 256
template <typename T>
static void diag_kernel(T * __restrict__ dst, const T * __restrict__ src,
const int64_t ne0, const int64_t ne1,
const int64_t ne2, const int64_t ne3,
const int64_t total_elements,
const sycl::nd_item<1> & item) {
const int64_t i = item.get_global_id(0);
if (i >= total_elements) {
return;
}
const int64_t i0 = i % ne0;
const int64_t i1 = (i / ne0) % ne1;
const int64_t i2 = (i / (ne0 * ne1)) % ne2;
const int64_t i3 = i / (ne0 * ne1 * ne2);
const int64_t dst_idx = ((i3 * ne2 + i2) * ne1 + i1) * ne0 + i0;
if (i0 == i1) {
const int64_t batch_idx = i3 * ne2 + i2;
dst[dst_idx] = src[batch_idx * ne0 + i0];
} else {
dst[dst_idx] = T(0);
}
(void)ne3;
}
inline void ggml_sycl_op_diag(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
const ggml_tensor * src0 = dst->src[0];
GGML_ASSERT(ggml_is_contiguous(dst));
GGML_ASSERT(ggml_is_contiguous(src0));
GGML_ASSERT(src0->ne[1] == 1);
dpct::queue_ptr stream = ctx.stream();
SYCL_CHECK(ggml_sycl_set_device(ctx.device));
const void * src0_d = src0->data;
void * dst_d = dst->data;
const int64_t ne0 = dst->ne[0];
const int64_t ne1 = dst->ne[1];
const int64_t ne2 = dst->ne[2];
const int64_t ne3 = dst->ne[3];
const int64_t n_elems = ggml_nelements(dst);
const int64_t num_blocks = (n_elems + SYCL_DIAG_BLOCK_SIZE - 1) / SYCL_DIAG_BLOCK_SIZE;
GGML_ASSERT(dst->type == GGML_TYPE_F32);
stream->parallel_for(
sycl::nd_range<1>(num_blocks * SYCL_DIAG_BLOCK_SIZE, SYCL_DIAG_BLOCK_SIZE),
[=](sycl::nd_item<1> item) {
diag_kernel(static_cast<float *>(dst_d),
static_cast<const float *>(src0_d),
ne0, ne1, ne2, ne3, n_elems, item);
});
}
void ggml_sycl_diag(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
ggml_sycl_op_diag(ctx, dst);
}

View file

@ -0,0 +1,5 @@
#pragma once
#include "common.hpp"
void ggml_sycl_diag(ggml_backend_sycl_context & ctx, ggml_tensor * dst);

View file

@ -0,0 +1,55 @@
#include "fill.hpp"
#include "common.hpp"
#define SYCL_FILL_BLOCK_SIZE 256
template <typename T>
static void fill_kernel(T * dst, const int64_t k, const T value,
const sycl::nd_item<1> & item) {
const int64_t i = (int64_t)item.get_global_id(0);
if (i >= k) {
return;
}
dst[i] = value;
}
inline void ggml_sycl_op_fill(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
GGML_ASSERT(ggml_is_contiguous(dst));
dpct::queue_ptr stream = ctx.stream();
SYCL_CHECK(ggml_sycl_set_device(ctx.device));
float value;
memcpy(&value, dst->op_params, sizeof(float));
const int64_t k = ggml_nelements(dst);
const int64_t num_blocks = (k + SYCL_FILL_BLOCK_SIZE - 1) / SYCL_FILL_BLOCK_SIZE;
void * dst_d = dst->data;
switch (dst->type) {
case GGML_TYPE_F32:
stream->parallel_for(
sycl::nd_range<1>(num_blocks * SYCL_FILL_BLOCK_SIZE, SYCL_FILL_BLOCK_SIZE),
[=](sycl::nd_item<1> item) {
fill_kernel(static_cast<float *>(dst_d), k, value, item);
});
break;
case GGML_TYPE_F16:
{
sycl::half h_value = sycl::half(value);
stream->parallel_for(
sycl::nd_range<1>(num_blocks * SYCL_FILL_BLOCK_SIZE, SYCL_FILL_BLOCK_SIZE),
[=](sycl::nd_item<1> item) {
fill_kernel(static_cast<sycl::half *>(dst_d), k, h_value, item);
});
}
break;
default:
GGML_ABORT("unsupported type");
}
}
void ggml_sycl_fill(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/0);
ggml_sycl_op_fill(ctx, dst);
}

View file

@ -0,0 +1,5 @@
#pragma once
#include "common.hpp"
void ggml_sycl_fill(ggml_backend_sycl_context & ctx, ggml_tensor * dst);

View file

@ -0,0 +1,172 @@
#include "solve_tri.hpp"
#include "common.hpp"
#include <oneapi/mkl/blas.hpp>
template <int n_template, int k_template>
static void solve_tri_f32_fast(const float * __restrict__ A,
const float * __restrict__ B,
float * __restrict__ X,
const int64_t ne02, [[maybe_unused]] const int64_t ne03,
const int64_t nb02, const int64_t nb03,
const int64_t nb12, const int64_t nb13,
const int64_t nb2, const int64_t nb3,
const int n_arg, const int k_arg,
const sycl::nd_item<2> & item, float * sA) {
const int n = n_template == 0 ? n_arg : n_template;
const int k = k_template == 0 ? k_arg : k_template;
const int batch_idx = item.get_group(1);
const int lane = item.get_local_id(1) % WARP_SIZE;
const int col_idx = item.get_local_id(0);
if (col_idx >= k) {
return;
}
const int64_t i03 = batch_idx / ne02;
const int64_t i02 = batch_idx % ne02;
const float * A_batch = (const float *) ((const char *) A + i02 * nb02 + i03 * nb03);
const float * B_batch = (const float *) ((const char *) B + i02 * nb12 + i03 * nb13);
float * X_batch = (float *) ((char *) X + i02 * nb2 + i03 * nb3);
const int offset = item.get_local_id(1) + item.get_local_id(0) * item.get_local_range(1);
#pragma unroll
for (int i = 0; i < n * n; i += k * WARP_SIZE) {
const int i0 = i + offset;
if (i0 < n * n) {
sA[i0] = A_batch[i0];
}
}
item.barrier(sycl::access::fence_space::local_space);
float x_low = (lane < n) ? B_batch[lane * k + col_idx] : 0.0f;
float x_high = (WARP_SIZE + lane < n) ? B_batch[(WARP_SIZE + lane) * k + col_idx] : 0.0f;
const int half = WARP_SIZE;
const int nrows_low = (n < half) ? n : half;
#pragma unroll
for (int row = 0; row < nrows_low; ++row) {
float sum = 0.0f;
if (lane < row) {
sum += sA[row * n + lane] * x_low;
}
sum = warp_reduce_sum<WARP_SIZE>(sum);
if (lane == row) {
x_low = (x_low - sum) / sA[row * n + row];
}
}
#pragma unroll
for (int row = half; row < n; ++row) {
float sum = sA[row * n + lane] * x_low;
const int j = half + lane;
if (j < row) {
sum += sA[row * n + j] * x_high;
}
sum = warp_reduce_sum<WARP_SIZE>(sum);
if (lane == row - half) {
x_high = (x_high - sum) / sA[row * n + row];
}
}
#pragma unroll
for (int rr = 0; rr < 2; ++rr) {
const int row = rr * WARP_SIZE + lane;
if (row < n) {
const float val = (row < half) ? x_low : x_high;
X_batch[row * k + col_idx] = val;
}
}
}
static void solve_tri_f32_mkl(dpct::queue_ptr stream,
const float * A, float * X,
int n, int k,
int64_t ne02, [[maybe_unused]] int64_t ne03,
int64_t nb02, [[maybe_unused]] int64_t nb03,
int64_t nb2, [[maybe_unused]] int64_t nb3) {
const float alpha = 1.0f;
const int64_t total_batches = ne02 * ne03;
if (total_batches == 0) {
return;
}
const int64_t stride_a = nb02 / sizeof(float);
const int64_t stride_x = nb2 / sizeof(float);
oneapi::mkl::blas::trsm_batch(
*stream,
oneapi::mkl::side::right,
oneapi::mkl::uplo::upper,
oneapi::mkl::transpose::nontrans,
oneapi::mkl::diag::nonunit,
k, n, alpha,
A, n, stride_a,
X, k, stride_x,
total_batches);
}
inline void ggml_sycl_op_solve_tri(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
const ggml_tensor * src0 = dst->src[0];
const ggml_tensor * src1 = dst->src[1];
GGML_ASSERT(ggml_is_contiguous(src0));
GGML_ASSERT(ggml_is_contiguous(src1));
GGML_ASSERT(src0->type == GGML_TYPE_F32);
dpct::queue_ptr stream = ctx.stream();
SYCL_CHECK(ggml_sycl_set_device(ctx.device));
const int n = src0->ne[0];
const int k = src1->ne[0];
const int64_t ne02 = src0->ne[2];
const int64_t ne03 = src0->ne[3];
GGML_ASSERT(n <= SYCL_SOLVE_TRI_MAX_N && k <= SYCL_SOLVE_TRI_MAX_K);
const float * A_d = static_cast<const float *>(src0->data);
const float * B_d = static_cast<const float *>(src1->data);
float * X_d = static_cast<float *>(dst->data);
if (X_d != B_d) {
const int64_t total_elements = (int64_t)n * k * ne02 * ne03;
stream->memcpy(X_d, B_d, total_elements * sizeof(float));
}
const int64_t nb02 = src0->nb[2];
const int64_t nb03 = src0->nb[3];
const int64_t nb12 = src1->nb[2];
const int64_t nb13 = src1->nb[3];
const int64_t nb2 = dst->nb[2];
const int64_t nb3 = dst->nb[3];
const int64_t total_batches = ne02 * ne03;
if (n <= 2 * WARP_SIZE && k <= 32) {
const int smem_size = 2 * WARP_SIZE * 2 * WARP_SIZE;
const sycl::range<2> grid(1, total_batches);
const sycl::range<2> block(k, WARP_SIZE);
stream->submit([&](sycl::handler & cgh) {
sycl::local_accessor<float, 1> smem_acc(sycl::range<1>(smem_size), cgh);
cgh.parallel_for(
sycl::nd_range<2>(grid * block, block),
[=](sycl::nd_item<2> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
solve_tri_f32_fast<0, 0>(A_d, B_d, X_d, ne02, ne03,
nb02, nb03, nb12, nb13, nb2, nb3,
n, k, item, get_pointer(smem_acc));
});
});
} else {
solve_tri_f32_mkl(stream, A_d, X_d, n, k, ne02, ne03, nb02, nb03, nb2, nb3);
}
}
void ggml_sycl_solve_tri(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
ggml_sycl_op_solve_tri(ctx, dst);
}

View file

@ -0,0 +1,8 @@
#pragma once
#include "common.hpp"
#define SYCL_SOLVE_TRI_MAX_N 64
#define SYCL_SOLVE_TRI_MAX_K 64
void ggml_sycl_solve_tri(ggml_backend_sycl_context & ctx, ggml_tensor * dst);

View file

@ -0,0 +1,156 @@
#include "ssm_scan.hpp"
#include "common.hpp"
template <int c_factor, int d_state>
static void ssm_scan_f32_group(
const float * __restrict__ src0, const float * __restrict__ src1, const float * __restrict__ src2,
const float * __restrict__ src3, const float * __restrict__ src4, const float * __restrict__ src5,
const int32_t * __restrict__ src6, float * __restrict__ dst,
const int src0_nb2, const int src0_nb3, const int src1_nb2, const int src1_nb3,
const int src2_nb1, const int src2_nb2, const int src3_nb1,
const int src4_nb2, const int src4_nb3, const int src5_nb2, const int src5_nb3,
const int64_t s_off, const int64_t n_head, const int64_t d_head, const int64_t n_group, const int64_t n_tok,
const sycl::nd_item<2> & item) {
const int lane = item.get_local_id(1) % WARP_SIZE;
const int warp = item.get_local_id(1) / WARP_SIZE;
const int warp_idx = item.get_group(1) * c_factor + warp;
const int seq_idx = item.get_group(0);
const int head_idx = warp_idx / d_head;
const int head_off = (warp_idx % d_head) * sizeof(float);
const int group_off = (head_idx / (n_head / n_group)) * d_state * sizeof(float);
const float * s0_warp = (const float *) ((const char *) src0 + src6[seq_idx] * src0_nb3 + head_idx * src0_nb2 + head_off * d_state);
const float * x_warp = (const float *) ((const char *) src1 + (seq_idx * src1_nb3) + (warp_idx * sizeof(float)));
const float * dt_warp = (const float *) ((const char *) src2 + (seq_idx * src2_nb2) + head_idx * sizeof(float));
const float * A_warp = (const float *) ((const char *) src3 + head_idx * src3_nb1);
const float * B_warp = (const float *) ((const char *) src4 + (seq_idx * src4_nb3) + (group_off));
const float * C_warp = (const float *) ((const char *) src5 + (seq_idx * src5_nb3) + (group_off));
float * y_warp = dst + (seq_idx * n_tok * n_head * d_head) + warp_idx;
float * s_warp = (float *) ((char *) dst + s_off + seq_idx * src0_nb3 + head_idx * src0_nb2 + head_off * d_state);
const int stride_x = src1_nb2 / sizeof(float);
const int stride_dt = src2_nb1 / sizeof(float);
const int stride_B = src4_nb2 / sizeof(float);
const int stride_C = src5_nb2 / sizeof(float);
const int stride_y = n_head * d_head;
float state[c_factor];
float state_sum = 0.0f;
#pragma unroll
for (int j = 0; j < c_factor; j++) {
state[j] = s0_warp[WARP_SIZE * j + lane];
}
for (int64_t i = 0; i < n_tok; i++) {
const float dt_val = dt_warp[i * stride_dt];
const float dt_soft_plus = (dt_val <= 20.0f ? sycl::log1p(sycl::exp(dt_val)) : dt_val);
state_sum = 0.0f;
const float dA = sycl::exp(dt_soft_plus * A_warp[0]);
const float x_dt = x_warp[i * stride_x] * dt_soft_plus;
#pragma unroll
for (int j = 0; j < c_factor; j++) {
const float B_val = B_warp[i * stride_B + WARP_SIZE * j + lane];
const float C_val = C_warp[i * stride_C + WARP_SIZE * j + lane];
state[j] = (state[j] * dA) + (B_val * x_dt);
state_sum += state[j] * C_val;
}
state_sum = warp_reduce_sum<WARP_SIZE>(state_sum);
if (lane == 0) {
y_warp[i * stride_y] = state_sum;
}
}
#pragma unroll
for (int j = 0; j < c_factor; j++) {
s_warp[WARP_SIZE * j + lane] = state[j];
}
}
static void ssm_scan_f32_sycl(
const float * src0, const float * src1, const float * src2, const float * src3,
const float * src4, const float * src5, const int32_t * src6, float * dst,
const int src0_nb2, const int src0_nb3, const int src1_nb2, const int src1_nb3, const int src2_nb1,
const int src2_nb2, const int src3_nb1, const int src4_nb2, const int src4_nb3, const int src5_nb2,
const int src5_nb3, const int64_t s_off, const int64_t d_state, const int64_t head_dim,
const int64_t n_head, const int64_t n_group, const int64_t n_tok, const int64_t n_seq,
dpct::queue_ptr stream) {
// NOTE: if you change conditions here, be sure to update the corresponding supports_op condition!
GGML_ASSERT(src3_nb1 == sizeof(float));
if (d_state == 128) {
constexpr int threads = 128;
constexpr int num_warps = threads / WARP_SIZE;
const sycl::range<2> grid(n_seq, (n_head * head_dim + num_warps - 1) / num_warps);
const sycl::range<2> block(1, threads);
stream->parallel_for(
sycl::nd_range<2>(grid * block, block),
[=](sycl::nd_item<2> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
ssm_scan_f32_group<128 / WARP_SIZE, 128>(
src0, src1, src2, src3, src4, src5, src6, dst,
src0_nb2, src0_nb3, src1_nb2, src1_nb3, src2_nb1, src2_nb2, src3_nb1,
src4_nb2, src4_nb3, src5_nb2, src5_nb3, s_off, n_head, head_dim, n_group, n_tok, item);
});
} else if (d_state == 256) {
constexpr int threads = 256;
constexpr int num_warps = threads / WARP_SIZE;
const sycl::range<2> grid(n_seq, (n_head * head_dim + num_warps - 1) / num_warps);
const sycl::range<2> block(1, threads);
stream->parallel_for(
sycl::nd_range<2>(grid * block, block),
[=](sycl::nd_item<2> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
ssm_scan_f32_group<256 / WARP_SIZE, 256>(
src0, src1, src2, src3, src4, src5, src6, dst,
src0_nb2, src0_nb3, src1_nb2, src1_nb3, src2_nb1, src2_nb2, src3_nb1,
src4_nb2, src4_nb3, src5_nb2, src5_nb3, s_off, n_head, head_dim, n_group, n_tok, item);
});
} else {
GGML_ABORT("ssm_scan: unsupported d_state (must be 128 or 256)");
}
}
inline void ggml_sycl_op_ssm_scan(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
const ggml_tensor * src0 = dst->src[0];
const ggml_tensor * src1 = dst->src[1];
const ggml_tensor * src2 = dst->src[2];
const ggml_tensor * src3 = dst->src[3];
const ggml_tensor * src4 = dst->src[4];
const ggml_tensor * src5 = dst->src[5];
const ggml_tensor * src6 = dst->src[6];
GGML_ASSERT(src0->type == GGML_TYPE_F32);
GGML_ASSERT(src6->type == GGML_TYPE_I32);
GGML_ASSERT(dst->type == GGML_TYPE_F32);
const int64_t nc = src0->ne[0];
const int64_t nr = src0->ne[1];
const int64_t nh = src1->ne[1];
const int64_t ng = src4->ne[1];
const int64_t n_t = src1->ne[2];
const int64_t n_s = src1->ne[3];
const int64_t s_off = ggml_nelements(src1) * sizeof(float);
GGML_ASSERT(ggml_nelements(src1) + nc * nr * nh * n_s == ggml_nelements(dst));
dpct::queue_ptr stream = ctx.stream();
SYCL_CHECK(ggml_sycl_set_device(ctx.device));
ssm_scan_f32_sycl(
static_cast<const float *>(src0->data), static_cast<const float *>(src1->data),
static_cast<const float *>(src2->data), static_cast<const float *>(src3->data),
static_cast<const float *>(src4->data), static_cast<const float *>(src5->data),
static_cast<const int32_t *>(src6->data), static_cast<float *>(dst->data),
src0->nb[2], src0->nb[3], src1->nb[2], src1->nb[3], src2->nb[1], src2->nb[2],
src3->nb[1], src4->nb[2], src4->nb[3], src5->nb[2], src5->nb[3],
s_off, nc, nr, nh, ng, n_t, n_s, stream);
}
void ggml_sycl_ssm_scan(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/7);
ggml_sycl_op_ssm_scan(ctx, dst);
}

View file

@ -0,0 +1,5 @@
#pragma once
#include "common.hpp"
void ggml_sycl_ssm_scan(ggml_backend_sycl_context & ctx, ggml_tensor * dst);

View file

@ -2154,11 +2154,11 @@ static void ggml_vk_create_pipeline_func(vk_device& device, vk_pipeline& pipelin
// Patch SPIR-V to enable RTE rounding for FP16, avoiding the need for
// separate shader variants compiled with -DRTE16.
std::vector<uint32_t> spv;
std::vector<uint32_t> spirv;
if (device->float_controls_rte_fp16) {
const uint32_t* spv_words = reinterpret_cast<const uint32_t *>(spv_data);
size_t word_count = spv_size / sizeof(uint32_t);
spv.assign(spv_words, spv_words + word_count);
spirv.assign(spv_words, spv_words + word_count);
// Find insertion points respecting SPIR-V layout order:
// Header(5) -> OpCapability -> OpExtension -> ... -> OpEntryPoint -> OpExecutionMode -> ...
@ -2168,9 +2168,9 @@ static void ggml_vk_create_pipeline_func(vk_device& device, vk_pipeline& pipelin
size_t exec_insert_pos = pos;
uint32_t entry_point_id = 0;
while (pos < spv.size()) {
uint32_t opcode = spv[pos] & spv::OpCodeMask;
uint32_t len = spv[pos] >> spv::WordCountShift;
while (pos < spirv.size()) {
uint32_t opcode = spirv[pos] & spv::OpCodeMask;
uint32_t len = spirv[pos] >> spv::WordCountShift;
if (len == 0) break;
if (opcode == spv::OpCapability) {
@ -2179,7 +2179,7 @@ static void ggml_vk_create_pipeline_func(vk_device& device, vk_pipeline& pipelin
} else if (opcode == spv::OpExtension) {
ext_insert_pos = pos + len;
} else if (opcode == spv::OpEntryPoint) {
entry_point_id = spv[pos + 2];
entry_point_id = spirv[pos + 2];
exec_insert_pos = pos + len;
} else if (opcode == spv::OpExecutionMode || opcode == spv::OpExecutionModeId) {
exec_insert_pos = pos + len;
@ -2194,7 +2194,7 @@ static void ggml_vk_create_pipeline_func(vk_device& device, vk_pipeline& pipelin
// OpExecutionMode %entrypoint RoundingModeRTE 16
uint32_t exec_mode[] = { (4u << spv::WordCountShift) | spv::OpExecutionMode, entry_point_id, spv::ExecutionModeRoundingModeRTE, 16 };
spv.insert(spv.begin() + exec_insert_pos, std::begin(exec_mode), std::end(exec_mode));
spirv.insert(spirv.begin() + exec_insert_pos, std::begin(exec_mode), std::end(exec_mode));
// OpExtension "SPV_KHR_float_controls"
const char ext_str[] = "SPV_KHR_float_controls";
@ -2202,13 +2202,13 @@ static void ggml_vk_create_pipeline_func(vk_device& device, vk_pipeline& pipelin
std::vector<uint32_t> extension(1 + ext_str_words, 0);
extension[0] = (uint32_t)((1 + ext_str_words) << spv::WordCountShift) | spv::OpExtension;
memcpy(&extension[1], ext_str, sizeof(ext_str));
spv.insert(spv.begin() + ext_insert_pos, extension.begin(), extension.end());
spirv.insert(spirv.begin() + ext_insert_pos, extension.begin(), extension.end());
// OpCapability RoundingModeRTE
uint32_t capability[] = { (2u << spv::WordCountShift) | spv::OpCapability, spv::CapabilityRoundingModeRTE };
spv.insert(spv.begin() + cap_insert_pos, std::begin(capability), std::end(capability));
spirv.insert(spirv.begin() + cap_insert_pos, std::begin(capability), std::end(capability));
shader_module_create_info = vk::ShaderModuleCreateInfo({}, spv.size() * sizeof(uint32_t), spv.data());
shader_module_create_info = vk::ShaderModuleCreateInfo({}, spirv.size() * sizeof(uint32_t), spirv.data());
}
pipeline->shader_module = device->device.createShaderModule(shader_module_create_info);

View file

@ -2443,6 +2443,8 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
MODEL_TENSOR.FFN_DOWN,
MODEL_TENSOR.FFN_UP,
MODEL_TENSOR.FFN_GATE_UP_EXP,
MODEL_TENSOR.FFN_GATE_EXP,
MODEL_TENSOR.FFN_UP_EXP,
MODEL_TENSOR.FFN_DOWN_EXP,
MODEL_TENSOR.ATTN_NORM,
MODEL_TENSOR.ATTN_POST_NORM,

View file

@ -2461,7 +2461,30 @@ public:
for (auto & [buft, mbuf] : mbufs_new) {
auto & mbuf_cur = mbufs[buft];
if (!mbuf_cur.buf || mbuf_cur.org.size() != mbuf.org.size() || mbuf_cur.total_size != mbuf.total_size) {
bool need_alloc = false;
need_alloc = need_alloc || (!mbuf_cur.buf);
need_alloc = need_alloc || (mbuf_cur.org.size() != mbuf.org.size());
need_alloc = need_alloc || (mbuf_cur.total_size != mbuf.total_size);
if (!need_alloc) {
for (size_t i = 0; i < mbuf_cur.org.size(); ++i) {
auto * org0 = mbuf_cur.org[i];
auto * org1 = mbuf.org[i];
if (!ggml_are_same_shape(org0, org1)) {
need_alloc = true;
break;
}
if (org0->view_src != org1->view_src || org0->view_offs != org1->view_offs) {
need_alloc = true;
break;
}
}
}
if (need_alloc) {
mbuf_cur = std::move(mbuf);
mbuf_cur.buf.reset(ggml_backend_alloc_ctx_tensors_from_buft(mbuf_cur.ctx.get(), buft));
@ -2525,6 +2548,31 @@ public:
mbufs_new[buft].total_size += rinfo.size;
}
for (auto & [buft, mbuf] : mbufs_new) {
ggml_init_params params = {
/*.mem_size =*/ mbuf.n_tensors*ggml_tensor_overhead(),
/*.mem_buffer =*/ NULL,
/*.no_alloc =*/ true,
};
mbuf.ctx.reset(ggml_init(params));
mbuf.org.reserve(mbuf.n_tensors);
}
for (const auto & rinfo : rinfos) {
auto * buft = ggml_backend_buffer_get_type(rinfo.tensor->buffer);
const int64_t n = rinfo.size/ggml_element_size(rinfo.tensor);
auto & mbuf = mbufs_new[buft];
mbuf.org.push_back(ggml_view_1d(mbuf.ctx.get(), rinfo.tensor, n, rinfo.offset));
auto & view = mbuf.org.back();
view->buffer = rinfo.tensor->buffer;
}
for (auto & [buft, mbuf] : mbufs_new) {
const auto & mbuf_cur = mbufs.at(buft);
@ -2533,9 +2581,11 @@ public:
}
for (size_t i = 0; i < mbuf_cur.org.size(); ++i) {
ggml_backend_tensor_copy(mbuf_cur.cpy[i], mbuf_cur.org[i]);
ggml_backend_tensor_copy(mbuf_cur.cpy[i], mbuf.org[i]);
}
}
GGML_ASSERT(buf_size == 0);
}
void read(void * dst, size_t size) override {

View file

@ -726,6 +726,10 @@ void llama_memory_recurrent::state_write(llama_io_write_i & io, llama_seq_id seq
cell_ranges.emplace_back(cell_range_begin, size);
}
if (flags % LLAMA_STATE_SEQ_FLAGS_ON_DEVICE && cell_ranges.size() > 1) {
GGML_ABORT("cannot save/load multiple ranges of cells to/from device memory\n");
}
// DEBUG CHECK: Sum of cell counts in ranges should equal the total cell count
uint32_t cell_count_check = 0;
for (const auto & range : cell_ranges) {

View file

@ -110,7 +110,13 @@ void llama_model_gemma4::load_arch_tensors(llama_model_loader &) {
layer.ffn_post_norm_2 = create_tensor(tn(LLM_TENSOR_FFN_POST_NORM_2, "weight", i), {n_embd}, 0);
// MoE FFN
layer.ffn_gate_up_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_UP_EXPS, "weight", i), {n_embd, n_ff_exp * 2, n_expert}, 0);
layer.ffn_gate_up_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_UP_EXPS, "weight", i), {n_embd, n_ff_exp * 2, n_expert}, TENSOR_NOT_REQUIRED);
if (layer.ffn_gate_up_exps == nullptr) {
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, 0);
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, 0);
}
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, 0);
// per-expert scale will be loaded as down_exps_s at the end of the current switch case
@ -286,8 +292,8 @@ llama_model_gemma4::graph::graph(const llama_model & model, const llm_graph_para
cur_moe = build_moe_ffn(cur_moe,
nullptr, // gate_inp
nullptr, // up_exps
nullptr, // gate_exps
model.layers[il].ffn_up_exps,
model.layers[il].ffn_gate_exps,
model.layers[il].ffn_down_exps,
nullptr, // exp_probs_b (not used for gemma4)
n_expert, n_expert_used,
@ -296,8 +302,8 @@ llama_model_gemma4::graph::graph(const llama_model & model, const llm_graph_para
LLAMA_EXPERT_GATING_FUNC_TYPE_SOFTMAX,
il, logits,
model.layers[il].ffn_gate_up_exps,
nullptr, // up_exps_s
nullptr, // gate_exps_s
model.layers[il].ffn_up_exps_s,
model.layers[il].ffn_gate_exps_s,
model.layers[il].ffn_down_exps_s);
cur_moe = build_norm(cur_moe,
model.layers[il].ffn_post_norm_2, nullptr,

File diff suppressed because one or more lines are too long

File diff suppressed because it is too large Load diff

View file

@ -3926,22 +3926,7 @@ void server_routes::init_routes() {
}},
{"object", "list"},
{"data", {
{
{"id", meta->model_name},
{"aliases", meta->model_aliases},
{"tags", meta->model_tags},
{"object", "model"},
{"created", std::time(0)},
{"owned_by", "llamacpp"},
{"meta", {
{"vocab_type", meta->model_vocab_type},
{"n_vocab", meta->model_vocab_n_tokens},
{"n_ctx_train", meta->model_n_ctx_train},
{"n_embd", meta->model_n_embd_inp},
{"n_params", meta->model_n_params},
{"size", meta->model_size},
}},
},
get_model_info(),
}}
};
@ -4155,6 +4140,26 @@ void server_routes::init_routes() {
};
}
json server_routes::get_model_info() const {
return json {
{"id", meta->model_name},
{"aliases", meta->model_aliases},
{"tags", meta->model_tags},
{"object", "model"},
{"created", std::time(0)},
{"owned_by", "llamacpp"},
{"meta", {
{"vocab_type", meta->model_vocab_type},
{"n_vocab", meta->model_vocab_n_tokens},
{"n_ctx", meta->slot_n_ctx},
{"n_ctx_train", meta->model_n_ctx_train},
{"n_embd", meta->model_n_embd_inp},
{"n_params", meta->model_n_params},
{"size", meta->model_size},
}},
};
}
std::unique_ptr<server_res_generator> server_routes::handle_slots_save(const server_http_req & req, int id_slot) {
auto res = create_response();
const json request_data = json::parse(req.body);

View file

@ -122,6 +122,10 @@ struct server_routes {
server_http_context::handler_t post_rerank;
server_http_context::handler_t get_lora_adapters;
server_http_context::handler_t post_lora_adapters;
// to be used in router mode
json get_model_info() const;
private:
std::unique_ptr<server_res_generator> handle_completions_impl(
const server_http_req & req,

View file

@ -4,7 +4,9 @@
#include <cpp-httplib/httplib.h>
#include <cstdlib>
#include <functional>
#include <future>
#include <string>
#include <thread>
@ -51,11 +53,51 @@ static void log_server_request(const httplib::Request & req, const httplib::Resp
SRV_DBG("response: %s\n", res.body.c_str());
}
// For Google Cloud Platform deployment compatibility
struct gcp_params {
bool enabled;
std::string path_health;
std::string path_predict;
int port;
// Ref: https://docs.cloud.google.com/vertex-ai/docs/predictions/custom-container-requirements#aip-variables
gcp_params() {
enabled = getenv("AIP_MODE", "") == "PREDICTION";
path_health = getenv("AIP_HEALTH_ROUTE", "", true); // default: using the route defined in server.cpp
path_predict = getenv("AIP_PREDICT_ROUTE", "/predict", true);
port = std::stoi(getenv("AIP_HTTP_PORT", "8080"));
}
static std::string getenv(const char * name, const std::string & default_value, bool ensure_leading_slash = false) {
const char * value = std::getenv(name);
if (value == nullptr || value[0] == '\0') {
return default_value;
}
std::string val = value;
if (ensure_leading_slash && !val.empty() && val[0] != '/') {
val.insert(val.begin(), '/');
}
return val;
}
};
bool server_http_context::init(const common_params & params) {
const gcp_params gcp;
path_prefix = params.api_prefix;
port = params.port;
hostname = params.hostname;
if (gcp.enabled) {
LOG_INF("%s: Google Cloud Platform compat: health route = %s, predict route = %s, port = %d\n", __func__, gcp.path_health.c_str(), gcp.path_predict.c_str(), gcp.port);
if (port != gcp.port) {
LOG_WRN("%s: Google Cloud Platform compat: overriding server port %d with AIP_HTTP_PORT %d\n", __func__, port, gcp.port);
}
port = gcp.port;
}
auto & srv = pimpl->srv;
#ifdef CPPHTTPLIB_OPENSSL_SUPPORT
@ -420,6 +462,7 @@ static void process_handler_response(server_http_req_ptr && request, server_http
}
void server_http_context::get(const std::string & path, const server_http_context::handler_t & handler) const {
handlers.emplace(path, handler);
pimpl->srv->Get(path_prefix + path, [handler](const httplib::Request & req, httplib::Response & res) {
server_http_req_ptr request = std::make_unique<server_http_req>(server_http_req{
get_params(req),
@ -436,6 +479,7 @@ void server_http_context::get(const std::string & path, const server_http_contex
}
void server_http_context::post(const std::string & path, const server_http_context::handler_t & handler) const {
handlers.emplace(path, handler);
pimpl->srv->Post(path_prefix + path, [handler](const httplib::Request & req, httplib::Response & res) {
std::string body = req.body;
std::map<std::string, uploaded_file> files;
@ -481,3 +525,176 @@ void server_http_context::post(const std::string & path, const server_http_conte
});
}
//
// Vertex AI Prediction protocol (AIP_PREDICT_ROUTE)
// https://cloud.google.com/vertex-ai/docs/predictions/custom-container-requirements
//
// Derives the camelCase @requestFormat alias for a registered path.
// e.g. "/v1/chat/completions" -> "chatCompletions", "/apply-template" -> "applyTemplate"
static std::string path_to_gcp_format(const std::string & path) {
std::string s = path;
if (s.size() > 3 && s[0] == '/' && s[1] == 'v' && s[2] == '1') {
s = s.substr(3);
}
if (!s.empty() && s[0] == '/') {
s = s.substr(1);
}
std::string result;
bool cap = false;
for (unsigned char c : s) {
if (c == ':') break; // stop before path parameters
if (c == '/' || c == '-' || c == '_') {
cap = true;
} else {
result += cap ? (char)std::toupper(c) : (char)c;
cap = false;
}
}
return result;
}
static json parse_gcp_predict_response(const server_http_res_ptr & res) {
if (res == nullptr) {
throw std::runtime_error("empty response from internal handler");
}
if (res->is_stream()) {
throw std::invalid_argument("predict route does not support streaming responses");
}
if (res->data.empty()) {
return nullptr;
}
try {
return json::parse(res->data);
} catch (...) {
return res->data;
}
}
void server_http_context::register_gcp_compat() {
const gcp_params gcp;
if (!gcp.enabled) {
// do nothing
return;
}
if (handlers.count(gcp.path_predict)) {
LOG_ERR("%s: AIP_PREDICT_ROUTE=%s conflicts with an existing llama-server route\n", __func__, gcp.path_predict.c_str());
exit(1);
}
// camelCase alias -> canonical path (first registration wins on collision)
// e.g. "chatCompletions" -> "/v1/chat/completions"
std::unordered_map<std::string, std::string> alias_to_path;
for (const auto & [path, _] : handlers) {
alias_to_path.emplace(path_to_gcp_format(path), path);
}
if (!gcp.path_health.empty()) {
auto health_handler = handlers.find("/health");
GGML_ASSERT(health_handler != handlers.end());
get(gcp.path_health, health_handler->second);
}
post(gcp.path_predict, [this, alias_to_path = std::move(alias_to_path)](const server_http_req & req) -> server_http_res_ptr {
static const auto build_error = [](const std::string & message, error_type type) -> json {
return json {{"error", format_error_response(message, type)}};
};
json data;
try {
data = json::parse(req.body);
} catch (const std::exception & e) {
auto res = std::make_unique<server_http_res>();
res->status = 400;
res->data = safe_json_to_str({{"error", format_error_response(e.what(), ERROR_TYPE_INVALID_REQUEST)}});
return res;
}
if (!data.is_object()) {
auto res = std::make_unique<server_http_res>();
res->status = 400;
res->data = safe_json_to_str({{"error", format_error_response("request body must be a JSON object", ERROR_TYPE_INVALID_REQUEST)}});
return res;
}
if (!data.contains("instances") || !data.at("instances").is_array()) {
auto res = std::make_unique<server_http_res>();
res->status = 400;
res->data = safe_json_to_str({{"error", format_error_response("request body must include an array field named instances", ERROR_TYPE_INVALID_REQUEST)}});
return res;
}
const json & instances = data.at("instances");
static const size_t MAX_INSTANCES = 128;
if (instances.size() > MAX_INSTANCES) {
auto res = std::make_unique<server_http_res>();
res->status = 400;
res->data = safe_json_to_str({{"error", format_error_response("instances array exceeds maximum size of " + std::to_string(MAX_INSTANCES), ERROR_TYPE_INVALID_REQUEST)}});
return res;
}
std::vector<std::future<json>> futures;
futures.reserve(instances.size());
for (const auto & instance : instances) {
futures.push_back(std::async(std::launch::async, [this, &req, &alias_to_path, instance]() -> json {
if (!instance.is_object()) {
return build_error("each instance must be a JSON object", ERROR_TYPE_INVALID_REQUEST);
}
if (!instance.contains("@requestFormat") || !instance.at("@requestFormat").is_string()) {
return build_error("each instance must include a string @requestFormat", ERROR_TYPE_INVALID_REQUEST);
}
try {
json payload = instance;
const std::string format = payload.at("@requestFormat").get<std::string>();
payload.erase("@requestFormat");
if (payload.contains("stream")) {
LOG_WRN("%s: ignoring client-provided stream field in instance, streaming is not supported in predict route\n", __func__);
payload["stream"] = false;
}
// accept both camelCase aliases (e.g. "chatCompletions") and direct paths
std::string dispatch_path;
auto it_alias = alias_to_path.find(format);
if (it_alias != alias_to_path.end()) {
dispatch_path = it_alias->second;
} else if (handlers.count(format)) {
dispatch_path = format;
} else {
return build_error("no handler registered for @requestFormat: " + format, ERROR_TYPE_INVALID_REQUEST);
}
const server_http_req internal_req {
req.params,
req.headers,
path_prefix + dispatch_path,
req.query_string,
payload.dump(),
{},
req.should_stop,
};
server_http_res_ptr internal_res = handlers.at(dispatch_path)(internal_req);
return parse_gcp_predict_response(internal_res);
} catch (const std::invalid_argument & e) {
return build_error(e.what(), ERROR_TYPE_INVALID_REQUEST);
} catch (const std::exception & e) {
return build_error(e.what(), ERROR_TYPE_SERVER);
} catch (...) {
return build_error("unknown error", ERROR_TYPE_SERVER);
}
}));
}
json predictions = json::array();
for (auto & future : futures) {
predictions.push_back(future.get());
}
auto res = std::make_unique<server_http_res>();
res->data = safe_json_to_str({{"predictions", predictions}});
return res;
});
}

View file

@ -67,6 +67,10 @@ struct server_http_context {
std::thread thread; // server thread
std::atomic<bool> is_ready = false;
// note: the handler should never throw exceptions
using handler_t = std::function<server_http_res_ptr(const server_http_req & req)>;
mutable std::unordered_map<std::string, handler_t> handlers;
std::string path_prefix;
std::string hostname;
int port;
@ -78,12 +82,13 @@ struct server_http_context {
bool start();
void stop() const;
// note: the handler should never throw exceptions
using handler_t = std::function<server_http_res_ptr(const server_http_req & req)>;
void get(const std::string & path, const handler_t & handler) const;
void post(const std::string & path, const handler_t & handler) const;
// Register the Google Cloud Platform (Vertex AI) compat (AIP_PREDICT_ROUTE env var, or /predict)
// Must be called AFTER all other API routes are registered
void register_gcp_compat();
// for debugging
std::string listening_address;
};

View file

@ -44,6 +44,7 @@ extern char **environ;
#define CMD_ROUTER_TO_CHILD_EXIT "cmd_router_to_child:exit"
#define CMD_CHILD_TO_ROUTER_READY "cmd_child_to_router:ready" // also sent when waking up from sleep
#define CMD_CHILD_TO_ROUTER_SLEEP "cmd_child_to_router:sleep"
#define CMD_CHILD_TO_ROUTER_INFO "cmd_child_to_router:info:" // followed by json string
// address for child process, this is needed because router may run on 0.0.0.0
// ref: https://github.com/ggml-org/llama.cpp/issues/17862
@ -718,10 +719,11 @@ void server_models::load(const std::string & name) {
// prepare new instance info
instance_t inst;
inst.meta = meta;
inst.meta.port = get_free_port();
inst.meta.status = SERVER_MODEL_STATUS_LOADING;
inst.meta.last_used = ggml_time_ms();
inst.meta = meta;
inst.meta.port = get_free_port();
inst.meta.status = SERVER_MODEL_STATUS_LOADING;
inst.meta.loaded_info = json{};
inst.meta.last_used = ggml_time_ms();
if (inst.meta.port <= 0) {
throw std::runtime_error("failed to get a port number");
@ -767,12 +769,14 @@ void server_models::load(const std::string & name) {
// read stdout/stderr and forward to main server log
// also handle status report from child process
if (stdout_file) {
char buffer[4096];
char buffer[128 * 1024]; // large buffer for storing info
while (fgets(buffer, sizeof(buffer), stdout_file) != nullptr) {
LOG("[%5d] %s", port, buffer);
std::string str(buffer);
if (string_starts_with(buffer, CMD_CHILD_TO_ROUTER_READY)) {
this->update_status(name, SERVER_MODEL_STATUS_LOADED, 0);
} else if (string_starts_with(buffer, CMD_CHILD_TO_ROUTER_INFO)) {
this->update_loaded_info(name, str);
} else if (string_starts_with(buffer, CMD_CHILD_TO_ROUTER_SLEEP)) {
this->update_status(name, SERVER_MODEL_STATUS_SLEEPING, 0);
}
@ -916,6 +920,29 @@ void server_models::update_status(const std::string & name, server_model_status
cv.notify_all();
}
void server_models::update_loaded_info(const std::string & name, std::string & raw_info) {
if (!string_starts_with(raw_info, CMD_CHILD_TO_ROUTER_INFO)) {
SRV_WRN("invalid loaded info format from child for model name=%s: %s\n", name.c_str(), raw_info.c_str());
return;
}
json info;
try {
info = json::parse(raw_info.substr(strlen(CMD_CHILD_TO_ROUTER_INFO)));
} catch (const std::exception & e) {
SRV_WRN("failed to parse loaded info from child for model name=%s: %s\n", name.c_str(), e.what());
return;
}
std::unique_lock<std::mutex> lk(mutex);
auto it = mapping.find(name);
if (it != mapping.end()) {
auto & meta = it->second.meta;
meta.loaded_info = info;
}
cv.notify_all();
}
void server_models::wait_until_loading_finished(const std::string & name) {
std::unique_lock<std::mutex> lk(mutex);
cv.wait(lk, [this, &name]() {
@ -994,12 +1021,14 @@ bool server_models::is_child_server() {
return router_port != nullptr;
}
std::thread server_models::setup_child_server(const std::function<void(int)> & shutdown_handler) {
std::thread server_models::setup_child_server(const std::function<void(int)> & shutdown_handler, const json & model_info) {
// send a notification to the router server that a model instance is ready
common_log_pause(common_log_main());
fflush(stdout);
fprintf(stdout, "%s\n", CMD_CHILD_TO_ROUTER_READY);
fflush(stdout);
fprintf(stdout, "%s%s\n", CMD_CHILD_TO_ROUTER_INFO, safe_json_to_str(model_info).c_str());
fflush(stdout);
common_log_resume(common_log_main());
// setup thread for monitoring stdin
@ -1176,7 +1205,8 @@ void server_models_routes::init_routes() {
status["exit_code"] = meta.exit_code;
status["failed"] = true;
}
models_json.push_back(json {
json model_info = json {
{"id", meta.name},
{"aliases", meta.aliases},
{"tags", meta.tags},
@ -1185,7 +1215,17 @@ void server_models_routes::init_routes() {
{"created", t}, // for OAI-compat
{"status", status},
// TODO: add other fields, may require reading GGUF metadata
});
};
// merge with loaded_info from the child process if available
if (meta.is_running()) {
for (auto it = meta.loaded_info.begin(); it != meta.loaded_info.end(); ++it) {
if (!model_info.contains(it.key())) {
model_info[it.key()] = it.value();
}
}
}
models_json.push_back(model_info);
}
res_ok(res, {
{"data", models_json},

View file

@ -63,6 +63,7 @@ struct server_model_meta {
server_model_status status = SERVER_MODEL_STATUS_UNLOADED;
int64_t last_used = 0; // for LRU unloading
std::vector<std::string> args; // args passed to the model instance, will be populated by render_args()
json loaded_info; // info to be reflected via /v1/models endpoint
int exit_code = 0; // exit code of the model instance process (only valid if status == FAILED)
int stop_timeout = 0; // seconds to wait before force-killing the model instance during shutdown
@ -145,6 +146,7 @@ public:
// update the status of a model instance (thread-safe)
void update_status(const std::string & name, server_model_status status, int exit_code);
void update_loaded_info(const std::string & name, std::string & raw_info);
// wait until the model instance is fully loaded (thread-safe)
// return when the model no longer in "loading" state
@ -163,7 +165,7 @@ public:
// notify the router server that a model instance is ready
// return the monitoring thread (to be joined by the caller)
static std::thread setup_child_server(const std::function<void(int)> & shutdown_handler);
static std::thread setup_child_server(const std::function<void(int)> & shutdown_handler, const json & model_info);
// notify the router server that the sleeping state has changed
static void notify_router_sleeping_state(bool sleeping);

View file

@ -204,6 +204,10 @@ int main(int argc, char ** argv) {
// Save & load slots
ctx_http.get ("/slots", ex_wrapper(routes.get_slots));
ctx_http.post("/slots/:id_slot", ex_wrapper(routes.post_slots));
// Google Cloud Platform (Vertex AI) compat
ctx_http.register_gcp_compat();
// CORS proxy (EXPERIMENTAL, only used by the Web UI for MCP)
if (params.webui_mcp_proxy) {
SRV_WRN("%s", "-----------------\n");
@ -334,7 +338,8 @@ int main(int argc, char ** argv) {
// optionally, notify router server that this instance is ready
std::thread monitor_thread;
if (server_models::is_child_server()) {
monitor_thread = server_models::setup_child_server(shutdown_handler);
json model_info = routes.get_model_info();
monitor_thread = server_models::setup_child_server(shutdown_handler, model_info);
}
// this call blocks the main thread until queue_tasks.terminate() is called

View file

@ -0,0 +1,60 @@
import pytest
from utils import *
server: ServerProcess
@pytest.fixture(autouse=True)
def create_server():
global server
server = ServerPreset.tinyllama2()
server.gcp_compat = True
def test_gcp_predict_camel_case():
global server
server.start()
res = server.make_request("POST", "/predict", data={
"instances": [
{
"@requestFormat": "chatCompletions",
"max_tokens": 8,
"messages": [
{"role": "user", "content": "What is the meaning of life?"},
],
}
],
})
assert res.status_code == 200
assert "predictions" in res.body
assert len(res.body["predictions"]) == 1
prediction = res.body["predictions"][0]
assert "choices" in prediction
assert len(prediction["choices"]) == 1
assert prediction["choices"][0]["message"]["role"] == "assistant"
assert len(prediction["choices"][0]["message"]["content"]) > 0
def test_gcp_predict_multiple_instances():
global server
server.n_slots = 2
server.start()
res = server.make_request("POST", "/predict", data={
"instances": [
{
"@requestFormat": "chatCompletions",
"max_tokens": 8,
"messages": [{"role": "user", "content": "Say hello"}],
},
{
"@requestFormat": "chatCompletions",
"max_tokens": 8,
"messages": [{"role": "user", "content": "Say world"}],
},
],
})
assert res.status_code == 200
assert len(res.body["predictions"]) == 2
for prediction in res.body["predictions"]:
assert "choices" in prediction
assert len(prediction["choices"][0]["message"]["content"]) > 0

View file

@ -108,6 +108,7 @@ class ServerProcess:
no_cache_idle_slots: bool = False
log_path: str | None = None
webui_mcp_proxy: bool = False
gcp_compat: bool = False
# session variables
process: subprocess.Popen | None = None
@ -122,6 +123,9 @@ class ServerProcess:
self.external_server = "DEBUG_EXTERNAL" in os.environ
def start(self, timeout_seconds: int = DEFAULT_HTTP_TIMEOUT) -> None:
env = {**os.environ}
if "LLAMA_CACHE" not in os.environ:
env["LLAMA_CACHE"] = "tmp"
if self.external_server:
print(f"[external_server]: Assuming external server running on {self.server_host}:{self.server_port}")
return
@ -248,6 +252,8 @@ class ServerProcess:
server_args.append("--no-cache-idle-slots")
if self.webui_mcp_proxy:
server_args.append("--webui-mcp-proxy")
if self.gcp_compat:
env["AIP_MODE"] = "PREDICTION"
args = [str(arg) for arg in [server_path, *server_args]]
print(f"tests: starting server with: {' '.join(args)}")
@ -268,7 +274,7 @@ class ServerProcess:
creationflags=flags,
stdout=self._log,
stderr=self._log if self._log != sys.stdout else sys.stdout,
env={**os.environ, "LLAMA_CACHE": "tmp"} if "LLAMA_CACHE" not in os.environ else None,
env=env,
)
server_instances.add(self)

View file

@ -2307,9 +2307,9 @@
}
},
"node_modules/@sveltejs/kit": {
"version": "2.57.1",
"resolved": "https://registry.npmjs.org/@sveltejs/kit/-/kit-2.57.1.tgz",
"integrity": "sha512-VRdSbB96cI1EnRh09CqmnQqP/YJvET5buj8S6k7CxaJqBJD4bw4fRKDjcarAj/eX9k2eHifQfDH8NtOh+ZxxPw==",
"version": "2.59.1",
"resolved": "https://registry.npmjs.org/@sveltejs/kit/-/kit-2.59.1.tgz",
"integrity": "sha512-d8OON70AphLdDesuTIl//M2O6fRTIicX8aYv8vhCiYEhTTI2OboKqey0Hu1A4VFhqwgqtq0vKDmPFGkw8kKmgw==",
"dev": true,
"license": "MIT",
"dependencies": {
@ -3640,9 +3640,9 @@
}
},
"node_modules/bits-ui": {
"version": "2.18.0",
"resolved": "https://registry.npmjs.org/bits-ui/-/bits-ui-2.18.0.tgz",
"integrity": "sha512-GLOBZRVy3hxNHIQ2MpD/+5aK9KcBFZRhUJtZ1UDABXdlVR4K6zFpgt4T+Rwuhf2sQzlc6yK1q/DprHPjwT4Pjw==",
"version": "2.18.1",
"resolved": "https://registry.npmjs.org/bits-ui/-/bits-ui-2.18.1.tgz",
"integrity": "sha512-KkemzKFH4T3gt3H+P86JcnAWExjByv/6vlwjm/BoCwTPHu03yiCdxbghdJLvFReQTe0acCAiRcKfmixxD6XvlA==",
"dev": true,
"license": "MIT",
"dependencies": {
@ -4856,9 +4856,9 @@
}
},
"node_modules/express-rate-limit": {
"version": "8.3.2",
"resolved": "https://registry.npmjs.org/express-rate-limit/-/express-rate-limit-8.3.2.tgz",
"integrity": "sha512-77VmFeJkO0/rvimEDuUC5H30oqUC4EyOhyGccfqoLebB0oiEYfM7nwPrsDsBL1gsTpwfzX8SFy2MT3TDyRq+bg==",
"version": "8.5.0",
"resolved": "https://registry.npmjs.org/express-rate-limit/-/express-rate-limit-8.5.0.tgz",
"integrity": "sha512-XKhFohWaSBdVJNTi5TaHziqnPkv04I9UQV6q1Wy7Ui6GGQZVW12ojDFwqer14EvCXxjvPG0CyWXx7cAXpALB4Q==",
"license": "MIT",
"dependencies": {
"ip-address": "10.1.0"
@ -7943,9 +7943,9 @@
}
},
"node_modules/postcss": {
"version": "8.5.6",
"resolved": "https://registry.npmjs.org/postcss/-/postcss-8.5.6.tgz",
"integrity": "sha512-3Ybi1tAuwAP9s0r1UQ2J4n5Y0G05bJkpUIO0/bI9MhwmD70S5aTWbXGBwxHrelT+XM1k6dM0pk+SwNkpTRN7Pg==",
"version": "8.5.14",
"resolved": "https://registry.npmjs.org/postcss/-/postcss-8.5.14.tgz",
"integrity": "sha512-SoSL4+OSEtR99LHFZQiJLkT59C5B1amGO1NzTwj7TT1qCUgUO6hxOvzkOYxD+vMrXBM3XJIKzokoERdqQq/Zmg==",
"dev": true,
"funding": [
{
@ -10084,9 +10084,9 @@
"license": "MIT"
},
"node_modules/uuid": {
"version": "13.0.0",
"resolved": "https://registry.npmjs.org/uuid/-/uuid-13.0.0.tgz",
"integrity": "sha512-XQegIaBTVUjSHliKqcnFqYypAd4S+WCYt5NIeRs6w/UAry7z8Y9j5ZwRRL4kzq9U3sD6v+85er9FvkEaBpji2w==",
"version": "13.0.2",
"resolved": "https://registry.npmjs.org/uuid/-/uuid-13.0.2.tgz",
"integrity": "sha512-vzi9uRZ926x4XV73S/4qQaTwPXM2JBj6/6lI/byHH1jOpCzb0zDbfytgA9LcN/hzb2l7WQSQnxITOVx5un/wGw==",
"dev": true,
"funding": [
"https://github.com/sponsors/broofa",
@ -10302,9 +10302,9 @@
}
},
"node_modules/vite-plugin-devtools-json/node_modules/uuid": {
"version": "11.1.0",
"resolved": "https://registry.npmjs.org/uuid/-/uuid-11.1.0.tgz",
"integrity": "sha512-0/A9rDy9P7cJ+8w1c9WD9V//9Wj15Ce2MPz8Ri6032usz+NfePxx5AcN3bN+r6ZL6jEo066/yNYB3tn4pQEx+A==",
"version": "11.1.1",
"resolved": "https://registry.npmjs.org/uuid/-/uuid-11.1.1.tgz",
"integrity": "sha512-vIYxrBCC/N/K+Js3qSN88go7kIfNPssr/hHCesKCQNAjmgvYS2oqr69kIufEG+O4+PfezOH4EbIeHCfFov8ZgQ==",
"dev": true,
"funding": [
"https://github.com/sponsors/broofa",

View file

@ -8,6 +8,7 @@
import { HealthCheckStatus } from '$lib/enums';
import type { MCPServerSettingsEntry } from '$lib/types';
import { goto } from '$app/navigation';
import { ROUTES } from '$lib/constants/routes';
interface Props {
onMcpSettingsClick?: () => void;
@ -52,7 +53,7 @@
function handleMcpSettingsClick() {
onMcpSettingsClick?.();
goto(`${hasMcpServers ? '' : '?add'}#/settings/mcp`);
goto(`${hasMcpServers ? '' : '?add'}${ROUTES.MCP_SERVERS}`);
}
</script>

View file

@ -12,6 +12,8 @@
import { useAttachmentMenu } from '$lib/hooks/use-attachment-menu.svelte';
import { AttachmentMenuItemId } from '$lib/enums';
import { PencilRuler } from '@lucide/svelte';
import { ROUTES, SETTINGS_SECTION_SLUGS } from '$lib/constants/routes';
import { RouterService } from '$lib/services/router.service';
interface Props {
class?: string;
@ -146,13 +148,16 @@
<div class="my-2 border-t"></div>
<a href="#/settings/mcp" class="flex items-center gap-3 px-3 py-2">
<a href={ROUTES.MCP_SERVERS} class="flex items-center gap-3 px-3 py-2">
<McpLogo class="inline h-4 w-4" />
<span class="text-sm">MCP Servers</span>
</a>
<a href="#/settings/chat/tools" class="flex items-center gap-3 px-3 py-2">
<a
href={RouterService.settings(SETTINGS_SECTION_SLUGS.TOOLS)}
class="flex items-center gap-3 px-3 py-2"
>
<PencilRuler class="inline h-4 w-4" />
<span class="text-sm">Tools</span>

View file

@ -13,6 +13,7 @@
import { conversationsStore } from '$lib/stores/conversations.svelte';
import { getFileTypeCategory } from '$lib/utils';
import { goto } from '$app/navigation';
import { ROUTES } from '$lib/constants/routes';
interface Props {
canSend?: boolean;
@ -100,7 +101,7 @@
{onSystemPromptClick}
{onMcpPromptClick}
{onMcpResourcesClick}
onMcpSettingsClick={() => goto('#/settings/mcp')}
onMcpSettingsClick={() => goto(ROUTES.MCP_SERVERS)}
/>
</div>
{/if}

View file

@ -17,6 +17,7 @@
import { parseFilesToMessageExtras } from '$lib/utils/browser-only';
import { deriveAgenticSections } from '$lib/utils';
import type { DatabaseMessageExtraMcpPrompt } from '$lib/types';
import { ROUTES } from '$lib/constants/routes';
interface Props {
class?: string;
@ -182,7 +183,7 @@
const conversationDeleted = await chatStore.removeSystemPromptPlaceholder(message.id);
if (conversationDeleted) {
goto(`#/`);
goto(ROUTES.START);
}
return;
@ -205,7 +206,7 @@
const conversationDeleted = await chatStore.removeSystemPromptPlaceholder(message.id);
if (conversationDeleted) {
goto(`#/`);
goto(ROUTES.START);
}
} else {
chatActions.delete(message);
@ -271,7 +272,7 @@
const conversationDeleted = await chatStore.removeSystemPromptPlaceholder(message.id);
isEditing = false;
if (conversationDeleted) {
goto(`#/`);
goto(ROUTES.START);
}
return;
}

View file

@ -85,10 +85,6 @@
editCtx.setUploadedFiles([...editCtx.editedUploadedFiles, ...processed]);
}
function handleUploadedFilesChange(files: ChatUploadedFile[]) {
editCtx.setUploadedFiles(files);
}
$effect(() => {
chatStore.setEditModeActive(handleFilesAdd);
@ -104,7 +100,7 @@
<ChatForm
value={editCtx.editedContent}
attachments={editCtx.editedExtras}
uploadedFiles={editCtx.editedUploadedFiles}
bind:uploadedFiles={editCtx.editedUploadedFiles}
placeholder="Edit your message..."
showMcpPromptButton
showAddButton={editCtx.messageRole === MessageRole.USER}
@ -112,7 +108,6 @@
onValueChange={editCtx.setContent}
onAttachmentRemove={handleAttachmentRemove}
onUploadedFileRemove={handleUploadedFileRemove}
onUploadedFilesChange={handleUploadedFilesChange}
onFilesAdd={handleFilesAdd}
onSubmit={handleSubmit}
/>

View file

@ -0,0 +1,83 @@
<script lang="ts">
import * as AlertDialog from '$lib/components/ui/alert-dialog';
import { Checkbox } from '$lib/components/ui/checkbox';
import Label from '$lib/components/ui/label/label.svelte';
import { Shield, ShieldOff } from '@lucide/svelte';
let {
open = $bindable(),
includeSensitiveData = $bindable(false),
onCancel,
onConfirm
}: {
open: boolean;
includeSensitiveData: boolean;
onCancel: () => void;
onConfirm: () => void;
} = $props();
function handleOpenChange(newOpen: boolean) {
if (!newOpen) {
onCancel();
}
}
</script>
<AlertDialog.Root {open} onOpenChange={handleOpenChange}>
<AlertDialog.Content>
<AlertDialog.Header>
<AlertDialog.Title class="flex items-center gap-2">
{#if includeSensitiveData}
<ShieldOff class="h-5 w-5 text-destructive" />
{:else}
<Shield class="h-5 w-5 text-destructive" />
{/if}
Export Settings
</AlertDialog.Title>
<AlertDialog.Description>
{#if includeSensitiveData}
<p class="text-amber-500">
Warning: This export will include sensitive data such as API keys and MCP server custom
headers (e.g., authorization tokens). Do not share this file with anyone you don't
trust.
</p>
{:else}
<p>
Sensitive data (API keys, MCP server custom headers) will not be included in the export
to protect your credentials.
</p>
{/if}
</AlertDialog.Description>
</AlertDialog.Header>
<div class="flex items-center gap-2 py-2">
<Checkbox id="include-sensitive" bind:checked={includeSensitiveData} />
<Label
for="include-sensitive"
class="text-sm leading-none peer-disabled:cursor-not-allowed peer-disabled:opacity-70"
>
{#if includeSensitiveData}
<span class="text-destructive">Include sensitive data (not recommended)</span>
{:else}
<span>Include sensitive data</span>
{/if}
</Label>
</div>
<AlertDialog.Footer>
<AlertDialog.Cancel onclick={onCancel}>Cancel</AlertDialog.Cancel>
<AlertDialog.Action
onclick={onConfirm}
class="bg-destructive text-white hover:bg-destructive/80"
>
{#if includeSensitiveData}
Export Anyway
{:else}
Export Without Sensitive Data
{/if}
</AlertDialog.Action>
</AlertDialog.Footer>
</AlertDialog.Content>
</AlertDialog.Root>

View file

@ -18,6 +18,37 @@
*/
export { default as DialogMcpServerAddNew } from './DialogMcpServerAddNew.svelte';
/**
* **DialogExportSettings** - Settings export dialog with sensitive data warning
*
* Dialog for exporting settings with an option to include or exclude
* sensitive data (API keys, MCP server custom headers). Defaults to excluding
* sensitive data for security. User must explicitly opt-in to include them.
*
* **Architecture:**
* - Uses ShadCN AlertDialog
* - Checkbox to toggle sensitive data inclusion (defaults to false)
* - Warning icon and message when sensitive data is included
* - Destructive variant for the action button when exporting with sensitive data
*
* **Features:**
* - Secure default: sensitive data excluded by default
* - User must explicitly opt-in to include sensitive data
* - Visual warning (ShieldOff icon) when sensitive data is included
* - Different action text based on sensitive data state
*
* @example
* ```svelte
* <DialogExportSettings
* bind:open={showExportSettings}
* bind:includeSensitiveData
* onConfirm={handleSettingsExport}
* onCancel={() => showExportSettings = false}
* />
* ```
*/
export { default as DialogExportSettings } from './DialogExportSettings.svelte';
/**
*
* CONFIRMATION DIALOGS

View file

@ -11,6 +11,8 @@
import ScrollArea from '$lib/components/ui/scroll-area/scroll-area.svelte';
import * as Sidebar from '$lib/components/ui/sidebar';
import Input from '$lib/components/ui/input/input.svelte';
import { ROUTES } from '$lib/constants/routes';
import { RouterService } from '$lib/services/router.service';
import {
conversationsStore,
conversations,
@ -159,7 +161,7 @@
}
handleMobileSidebarItemClick();
await goto(`#/chat/${id}`);
await goto(RouterService.chat(id));
}
function handleStopGeneration(id: string) {
@ -171,7 +173,7 @@
<ScrollArea class="h-full flex-1">
<Sidebar.Header class="gap-4 bg-sidebar/50 p-3 backdrop-blur-lg md:pt-4 md:pb-2">
<div class="flex items-center justify-between">
<a href="#/" onclick={handleMobileSidebarItemClick}>
<a href={ROUTES.START} onclick={handleMobileSidebarItemClick}>
<h1 class="inline-flex items-center gap-1 px-2 text-xl font-semibold">{APP_NAME}</h1>
</a>

View file

@ -11,8 +11,10 @@
import { DropdownMenuActions } from '$lib/components/app';
import * as Tooltip from '$lib/components/ui/tooltip';
import { FORK_TREE_DEPTH_PADDING } from '$lib/constants';
import { RouterService } from '$lib/services/router.service';
import { getAllLoadingChats } from '$lib/stores/chat.svelte';
import { conversationsStore } from '$lib/stores/conversations.svelte';
import { TruncatedText } from '$lib/components/app';
import { onMount } from 'svelte';
interface Props {
@ -112,7 +114,7 @@
<Tooltip.Root>
<Tooltip.Trigger>
<a
href="#/chat/{conversation.forkedFromConversationId}"
href={RouterService.chat(conversation.forkedFromConversationId)}
class="flex shrink-0 items-center text-muted-foreground transition-colors hover:text-foreground"
>
<GitBranch class="h-3.5 w-3.5" />
@ -148,9 +150,7 @@
</Tooltip.Root>
{/if}
<span class="truncate text-sm font-medium">
{conversation.name}
</span>
<TruncatedText text={conversation.name} class="text-sm font-medium" showTooltip={false} />
</div>
{#if renderActionsDropdown}

View file

@ -7,6 +7,8 @@
import Label from '$lib/components/ui/label/label.svelte';
import { serverStore, serverLoading } from '$lib/stores/server.svelte';
import { config, settingsStore } from '$lib/stores/settings.svelte';
import { SETTINGS_KEYS } from '$lib/constants';
import { ROUTES } from '$lib/constants/routes';
import { fade, fly, scale } from 'svelte/transition';
import { KeyboardKey } from '$lib/enums';
@ -63,7 +65,7 @@
try {
// Update the API key in settings first
settingsStore.updateConfig('apiKey', apiKeyInput.trim());
settingsStore.updateConfig(SETTINGS_KEYS.API_KEY, apiKeyInput.trim());
// Test the API key by making a real request to the server
const response = await fetch(`${base}/props`, {
@ -79,7 +81,7 @@
// Show success state briefly, then navigate to home
setTimeout(() => {
goto(`#/`);
goto(ROUTES.START);
}, 1000);
} else {
// API key is invalid - User Story A

View file

@ -1,11 +1,11 @@
<script lang="ts">
import SettingsChatFooter from './SettingsChatFooter.svelte';
import SettingsChatFields from './SettingsChatFields.svelte';
import SettingsChatToolsTab from './SettingsChatToolsTab.svelte';
import SettingsChatImportExportTab from './SettingsChatImportExportTab.svelte';
import {
SettingsChatDesktopSidebar,
SettingsChatMobileHeader
SettingsChatFields,
SettingsChatImportExportTab,
SettingsChatMobileHeader,
SettingsChatToolsTab,
SettingsFooter
} from '$lib/components/app/settings';
import { config, settingsStore } from '$lib/stores/settings.svelte';
import {
@ -15,6 +15,7 @@
SETTINGS_SECTION_TITLES,
type SettingsSection
} from '$lib/constants';
import { RouterService } from '$lib/services/router.service';
import { setMode } from 'mode-watcher';
import { ColorMode } from '$lib/enums/ui';
import { fade } from 'svelte/transition';
@ -22,7 +23,8 @@
import { page } from '$app/state';
import { setChatSettingsConfigContext } from '$lib/contexts';
import { settingsReferrer } from '$lib/stores/settings-referrer.svelte';
import { modelsStore } from '$lib/stores/models.svelte';
import { isRouterMode } from '$lib/stores/server.svelte';
interface Props {
initialSection?: string;
getSectionHref?: (section: SettingsSection) => string;
@ -33,14 +35,30 @@
let activeSlug = $derived(
initialSection ?? (page.params as Record<string, string | undefined>).section ?? 'general'
);
let currentSection = $derived(
SETTINGS_CHAT_SECTIONS.find((section) => section.slug === activeSlug) ||
SETTINGS_CHAT_SECTIONS[0]
);
let localConfig: SettingsConfigType = $state({ ...config() });
let mobileHeader: { updateCarousel: () => void } | undefined;
let fetchInitiated = false;
$effect(() => {
if (isRouterMode() && currentSection.fields && !fetchInitiated) {
fetchInitiated = true;
void modelsStore
.fetch()
.then(() => modelsStore.fetchRouterModels())
.then(() => modelsStore.fetchModalitiesForLoadedModels())
.then(() => modelsStore.ensureFirstModelSelected());
}
});
function handleThemeChange(newTheme: string) {
localConfig.theme = newTheme;
setMode(newTheme as ColorMode);
@ -110,13 +128,15 @@
<SettingsChatDesktopSidebar
sections={SETTINGS_CHAT_SECTIONS}
isActive={(section: SettingsSection) => section.slug === activeSlug}
getHref={getSectionHref ?? ((section: SettingsSection) => `#/settings/chat/${section.slug}`)}
getHref={getSectionHref ??
((section: SettingsSection) => RouterService.settings(section.slug))}
/>
<SettingsChatMobileHeader
sections={SETTINGS_CHAT_SECTIONS}
isActive={(section: SettingsSection) => section.slug === activeSlug}
getHref={getSectionHref ?? ((section: SettingsSection) => `#/settings/chat/${section.slug}`)}
getHref={getSectionHref ??
((section: SettingsSection) => RouterService.settings(section.slug))}
bind:this={mobileHeader}
/>
@ -149,7 +169,7 @@
</div>
</div>
<SettingsChatFooter onReset={handleReset} onSave={handleSave} />
<SettingsFooter onReset={handleReset} onSave={handleSave} />
</div>
</div>
</div>

View file

@ -9,9 +9,9 @@
import { SettingsFieldType } from '$lib/enums/settings';
import { settingsStore } from '$lib/stores/settings.svelte';
import { serverStore } from '$lib/stores/server.svelte';
import { modelsStore, selectedModelName } from '$lib/stores/models.svelte';
import { modelsStore, selectedModelName, propsCacheVersion } from '$lib/stores/models.svelte';
import { normalizeFloatingPoint } from '$lib/utils/precision';
import SettingsChatParameterSourceIndicator from './SettingsChatParameterSourceIndicator.svelte';
import { SettingsChatParameterSourceIndicator } from '$lib/components/app/settings';
import type { Component } from 'svelte';
interface Props {
@ -23,13 +23,19 @@
let { fields, localConfig, onConfigChange, onThemeChange }: Props = $props();
// server sampling defaults for placeholders
let sp = $derived.by(() => {
let currentModelParams = $derived.by(() => {
propsCacheVersion();
if (serverStore.isRouterMode) {
const m = selectedModelName();
if (m) {
const p = modelsStore.getModelProps(m);
return (p?.default_generation_settings?.params ?? {}) as Record<string, unknown>;
const currentModelName = selectedModelName();
if (currentModelName) {
const currentModelProps = modelsStore.getModelProps(currentModelName);
return (currentModelProps?.default_generation_settings?.params ?? {}) as Record<
string,
unknown
>;
}
}
return (serverStore.defaultParams ?? {}) as Record<string, unknown>;
@ -40,7 +46,7 @@
<div class="space-y-2">
{#if field.type === SettingsFieldType.INPUT}
{@const currentValue = String(localConfig[field.key] ?? '')}
{@const serverDefault = sp[field.key]}
{@const serverDefault = currentModelParams[field.key]}
{@const isCustomRealTime = (() => {
if (serverDefault == null) return false;
if (currentValue === '') return false;
@ -78,8 +84,8 @@
// Update local config immediately for real-time badge feedback
onConfigChange(field.key, e.currentTarget.value);
}}
placeholder={sp[field.key] != null
? `Default: ${normalizeFloatingPoint(sp[field.key])}`
placeholder={currentModelParams[field.key] != null
? `Default: ${normalizeFloatingPoint(currentModelParams[field.key])}`
: ''}
class="w-full {isCustomRealTime ? 'pr-8' : ''}"
/>
@ -104,13 +110,15 @@
</p>
{/if}
{:else if field.type === SettingsFieldType.TEXTAREA}
<Label for={field.key} class="block flex items-center gap-1.5 text-sm font-medium">
{field.label}
{#if field.label}
<Label for={field.key} class="block flex items-center gap-1.5 text-sm font-medium">
{field.label}
{#if field.isExperimental}
<FlaskConical class="h-3.5 w-3.5 text-muted-foreground" />
{/if}
</Label>
{#if field.isExperimental}
<FlaskConical class="h-3.5 w-3.5 text-muted-foreground" />
{/if}
</Label>
{/if}
<Textarea
id={field.key}
@ -131,7 +139,8 @@
<Checkbox
id="showSystemMessage"
checked={Boolean(localConfig.showSystemMessage ?? true)}
onCheckedChange={(checked) => onConfigChange('showSystemMessage', Boolean(checked))}
onCheckedChange={(checked) =>
onConfigChange(SETTINGS_KEYS.SHOW_SYSTEM_MESSAGE, Boolean(checked))}
/>
<Label for="showSystemMessage" class="cursor-pointer text-sm font-normal">
@ -145,7 +154,7 @@
opt.value === localConfig[field.key]
)}
{@const currentValue = localConfig[field.key]}
{@const serverDefault = sp[field.key]}
{@const serverDefault = currentModelParams[field.key]}
{@const isCustomRealTime = (() => {
if (serverDefault == null) return false;
if (currentValue === '' || currentValue === undefined) return false;

View file

@ -0,0 +1,62 @@
<script lang="ts">
import type { Component } from 'svelte';
import { Button, type ButtonVariant } from '$lib/components/ui/button';
let {
title,
description,
IconComponent,
buttonText,
onclick,
titleClass,
buttonVariant,
buttonClass,
wrapperClass,
summary
}: {
title: string;
description: string;
IconComponent: Component;
buttonText: string;
onclick: () => void;
titleClass?: string;
buttonVariant?: ButtonVariant;
buttonClass?: string;
wrapperClass?: string;
summary?: { show: boolean; verb: string; items: DatabaseConversation[] };
} = $props();
let sectionButtonClass = $derived(buttonClass ?? 'justify-start justify-self-start md:w-auto');
let sectionButtonVariant = $derived(buttonVariant ?? 'outline');
</script>
<div class="grid gap-1 {wrapperClass ?? ''}">
<h4 class="mt-0 mb-2 text-sm font-medium {titleClass ?? ''}">{title}</h4>
<p class="mb-4 text-sm text-muted-foreground">{description}</p>
<Button class={sectionButtonClass} {onclick} variant={sectionButtonVariant}>
<IconComponent class="mr-2 h-4 w-4" />
{buttonText}
</Button>
{#if summary && summary.show && summary.items.length > 0}
<div class="mt-4 grid overflow-x-auto rounded-lg border border-border/50 bg-muted/30 p-4">
<h5 class="mb-2 text-sm font-medium">
{summary.verb}
{summary.items.length} conversation{summary.items.length === 1 ? '' : 's'}
</h5>
<ul class="space-y-1 text-sm text-muted-foreground">
{#each summary.items.slice(0, 10) as conv (conv.id)}
<li class="truncate">{conv.name || 'Untitled conversation'}</li>
{/each}
{#if summary.items.length > 10}
<li class="italic">... and {summary.items.length - 10} more</li>
{/if}
</ul>
</div>
{/if}
</div>

View file

@ -1,21 +1,18 @@
<script lang="ts">
import type { Component } from 'svelte';
import { Download, Upload, Trash2 } from '@lucide/svelte';
import { Button, type ButtonVariant } from '$lib/components/ui/button';
import { DialogConversationSelection, DialogConfirmation } from '$lib/components/app';
import {
DialogConversationSelection,
DialogConfirmation,
DialogExportSettings
} from '$lib/components/app';
import { createMessageCountMap } from '$lib/utils';
import { settingsStore } from '$lib/stores/settings.svelte';
import { conversationsStore, conversations } from '$lib/stores/conversations.svelte';
import { toast } from 'svelte-sonner';
import { fade } from 'svelte/transition';
import { ConversationSelectionMode, HtmlInputType, FileExtensionText } from '$lib/enums';
interface SectionOpts {
wrapperClass?: string;
titleClass?: string;
buttonVariant?: ButtonVariant;
buttonClass?: string;
summary?: { show: boolean; verb: string; items: DatabaseConversation[] };
}
import SettingsChatImportExportSection from './SettingsChatImportExportSection.svelte';
import SettingsGroup from '$lib/components/app/settings/SettingsGroup.svelte';
let exportedConversations = $state<DatabaseConversation[]>([]);
let importedConversations = $state<DatabaseConversation[]>([]);
@ -33,6 +30,82 @@
// Delete functionality state
let showDeleteDialog = $state(false);
// Settings import/export state
let showSettingsExportSummary = $state(false);
let showSettingsImportSummary = $state(false);
let showSettingsExportDialog = $state(false);
let includeSensitiveData = $state(false);
function handleSettingsExport() {
showSettingsExportDialog = true;
includeSensitiveData = false;
}
function handleSettingsExportConfirm() {
showSettingsExportDialog = false;
try {
const data = settingsStore.exportSettings(includeSensitiveData);
const blob = new Blob([JSON.stringify(data, null, 2)], { type: 'application/json' });
const url = URL.createObjectURL(blob);
const a = document.createElement('a');
a.href = url;
a.download = `llama_settings_${new Date().toISOString().split('T')[0]}.json`;
document.body.appendChild(a);
a.click();
document.body.removeChild(a);
URL.revokeObjectURL(url);
showSettingsExportSummary = true;
showSettingsImportSummary = false;
toast.success('Settings exported');
} catch (err) {
console.error('Failed to export settings:', err);
toast.error('Failed to export settings');
}
}
function handleSettingsExportCancel() {
showSettingsExportDialog = false;
}
function handleSettingsImport() {
try {
const input = document.createElement('input');
input.type = HtmlInputType.FILE;
input.accept = FileExtensionText.JSON;
input.onchange = async (e) => {
const file = (e.target as HTMLInputElement)?.files?.[0];
if (!file) return;
try {
const text = await file.text();
const data = JSON.parse(text);
if (!data || typeof data !== 'object' || !data.config) {
toast.error('Invalid settings file: missing config');
return;
}
settingsStore.importSettings(data);
showSettingsImportSummary = true;
showSettingsExportSummary = false;
toast.success('Settings imported successfully');
} catch (err) {
console.error('Failed to import settings:', err);
toast.error('Failed to import settings');
}
};
input.click();
} catch (err) {
console.error('Failed to open file picker:', err);
toast.error('Failed to open file picker');
}
}
async function handleExportClick() {
try {
const allConversations = conversations();
@ -181,94 +254,66 @@
}
</script>
{#snippet summaryList(show: boolean, verb: string, items: DatabaseConversation[])}
{#if show && items.length > 0}
<div class="mt-4 grid overflow-x-auto rounded-lg border border-border/50 bg-muted/30 p-4">
<h5 class="mb-2 text-sm font-medium">
{verb}
{items.length} conversation{items.length === 1 ? '' : 's'}
</h5>
<div class="space-y-12" in:fade={{ duration: 150 }}>
<SettingsGroup title="Conversations">
<SettingsChatImportExportSection
title="Export"
description="Download your conversations as a JSON file. This includes all messages, attachments, and conversation history."
IconComponent={Download}
buttonText="Export conversations"
onclick={handleExportClick}
summary={{ show: showExportSummary, verb: 'Exported', items: exportedConversations }}
/>
<ul class="space-y-1 text-sm text-muted-foreground">
{#each items.slice(0, 10) as conv (conv.id)}
<li class="truncate">{conv.name || 'Untitled conversation'}</li>
{/each}
<SettingsChatImportExportSection
title="Import"
description="Import one or more conversations from a previously exported JSON file. This will merge with your existing conversations."
IconComponent={Upload}
buttonText="Import conversations"
onclick={handleImportClick}
summary={{ show: showImportSummary, verb: 'Imported', items: importedConversations }}
/>
{#if items.length > 10}
<li class="italic">... and {items.length - 10} more</li>
{/if}
</ul>
</div>
{/if}
{/snippet}
<SettingsChatImportExportSection
title="Delete All"
description="Permanently delete all conversations and their messages. This action cannot be undone. Consider exporting your conversations first if you want to keep a backup."
IconComponent={Trash2}
buttonText="Delete all conversations"
onclick={handleDeleteAllClick}
titleClass="text-destructive"
buttonVariant="destructive"
buttonClass="text-destructive-foreground justify-start justify-self-start bg-destructive hover:bg-destructive/80 md:w-auto"
/>
</SettingsGroup>
{#snippet section(
title: string,
description: string,
IconComponent: Component,
buttonText: string,
onclick: () => void,
opts: SectionOpts
)}
{@const buttonClass = opts?.buttonClass ?? 'justify-start justify-self-start md:w-auto'}
{@const buttonVariant = opts?.buttonVariant ?? 'outline'}
<div class="grid gap-1 {opts?.wrapperClass ?? ''}">
<h4 class="mt-0 mb-2 text-sm font-medium {opts?.titleClass ?? ''}">{title}</h4>
<SettingsGroup title="Settings">
<SettingsChatImportExportSection
title="Export"
description="Export your chat settings and preferences as a JSON file."
IconComponent={Download}
buttonText="Export settings"
onclick={handleSettingsExport}
summary={{ show: showSettingsExportSummary, verb: 'Exported', items: [] }}
/>
<p class="mb-4 text-sm text-muted-foreground">{description}</p>
<Button class={buttonClass} {onclick} variant={buttonVariant}>
<IconComponent class="mr-2 h-4 w-4" />
{buttonText}
</Button>
{#if opts?.summary}
{@render summaryList(opts.summary.show, opts.summary.verb, opts.summary.items)}
{/if}
</div>
{/snippet}
<div class="space-y-6" in:fade={{ duration: 150 }}>
<div class="space-y-6">
{@render section(
'Export Conversations',
'Download all your conversations as a JSON file. This includes all messages, attachments, and conversation history.',
Download,
'Export conversations',
handleExportClick,
{ summary: { show: showExportSummary, verb: 'Exported', items: exportedConversations } }
)}
{@render section(
'Import Conversations',
'Import one or more conversations from a previously exported JSON file. This will merge with your existing conversations.',
Upload,
'Import conversations',
handleImportClick,
{
wrapperClass: 'border-t border-border/30 pt-6',
summary: { show: showImportSummary, verb: 'Imported', items: importedConversations }
}
)}
{@render section(
'Delete All Conversations',
'Permanently delete all conversations and their messages. This action cannot be undone. Consider exporting your conversations first if you want to keep a backup.',
Trash2,
'Delete all conversations',
handleDeleteAllClick,
{
wrapperClass: 'border-t border-border/30 pt-4',
titleClass: 'text-destructive',
buttonVariant: 'destructive',
buttonClass:
'text-destructive-foreground justify-start justify-self-start bg-destructive hover:bg-destructive/80 md:w-auto'
}
)}
</div>
<SettingsChatImportExportSection
title="Import"
description="Import chat settings from a previously exported JSON file. This will merge with your existing settings."
IconComponent={Upload}
buttonText="Import settings"
onclick={handleSettingsImport}
summary={{ show: showSettingsImportSummary, verb: 'Imported', items: [] }}
/>
</SettingsGroup>
</div>
<DialogExportSettings
bind:open={showSettingsExportDialog}
bind:includeSensitiveData
onConfirm={handleSettingsExportConfirm}
onCancel={handleSettingsExportCancel}
/>
<DialogConversationSelection
conversations={availableConversations}
{messageCountMap}

View file

@ -0,0 +1,18 @@
<script lang="ts">
import type { Snippet } from 'svelte';
interface Props {
title: string;
children: Snippet;
}
let { title, children }: Props = $props();
</script>
<div>
<h3 class="mb-6 text-base font-semibold">{title}</h3>
<div class="space-y-8">
{@render children()}
</div>
</div>

View file

@ -19,32 +19,6 @@ export { default as SettingsChatDesktopSidebar } from './SettingsChatDesktopSide
*/
export { default as SettingsChatMobileHeader } from './SettingsChatMobileHeader.svelte';
/**
* Settings Import/Export panel.
* Provides UI for importing and exporting chat conversations.
*/
export { default as SettingsChatImportExportTab } from './SettingsChat/SettingsChatImportExportTab.svelte';
/**
* MCP Servers configuration panel.
* Provides UI for managing Model Context Protocol (MCP) server connections.
*/
export { default as SettingsMcpServers } from './SettingsMcpServers.svelte';
/**
* Footer with save/cancel buttons for settings panel. Positioned at bottom
* of settings dialog. Save button commits form state to config store,
* cancel button triggers reset and close.
*/
export { default as SettingsChatFooter } from './SettingsChat/SettingsChatFooter.svelte';
/**
* Form fields renderer for individual settings. Generates appropriate input
* components based on field type (text, number, select, checkbox, textarea).
* Handles validation, help text display, and parameter source indicators.
*/
export { default as SettingsChatFields } from './SettingsChat/SettingsChatFields.svelte';
/**
* Badge indicating parameter source for sampling settings. Shows one of:
* - **Custom**: User has explicitly set this value (orange badge)
@ -54,6 +28,44 @@ export { default as SettingsChatFields } from './SettingsChat/SettingsChatFields
*/
export { default as SettingsChatParameterSourceIndicator } from './SettingsChat/SettingsChatParameterSourceIndicator.svelte';
/**
* Section wrapper for settings panels. Displays a title heading with
* child content in a structured layout.
*/
export { default as SettingsGroup } from './SettingsGroup.svelte';
/**
* Footer with save/cancel buttons for settings panel. Positioned at bottom
* of settings dialog. Save button commits form state to config store,
* cancel button triggers reset and close.
*/
export { default as SettingsFooter } from './SettingsFooter.svelte';
/**
* Settings Import/Export panel.
* Provides UI for importing and exporting chat conversations.
*/
export { default as SettingsChatImportExportTab } from './SettingsChat/SettingsChatImportExportTab.svelte';
/**
* Section wrapper for import/export sections. Displays a title, description,
* icon button, and optional summary of recent actions.
*/
export { default as SettingsChatImportExportSection } from './SettingsChat/SettingsChatImportExportSection.svelte';
/**
* MCP Servers configuration panel.
* Provides UI for managing Model Context Protocol (MCP) server connections.
*/
export { default as SettingsMcpServers } from './SettingsMcpServers.svelte';
/**
* Form fields renderer for individual settings. Generates appropriate input
* components based on field type (text, number, select, checkbox, textarea).
* Handles validation, help text display, and parameter source indicators.
*/
export { default as SettingsChatFields } from './SettingsChat/SettingsChatFields.svelte';
/**
* **SettingsChatToolsTab** - Tools configuration tab for chat settings
*

View file

@ -29,12 +29,12 @@ export * from './message-export';
export * from './model-id';
export * from './precision';
export * from './processing-info';
export * from './settings-config';
export * from './settings-fields';
export * from './routes';
export * from './settings-keys';
export * from './settings-sections';
export * from './settings-registry';
export * from './supported-file-types';
export * from './table-html-restorer';
export * from './title-generation';
export * from './tools';
export * from './tooltip-config';
export * from './ui';

View file

@ -0,0 +1,26 @@
export const NEW_CHAT_PARAM = 'new_chat';
/** Settings section slugs — used for routes and navigation. */
export const SETTINGS_SECTION_SLUGS = {
GENERAL: 'general',
DISPLAY: 'display',
SAMPLING: 'sampling',
PENALTIES: 'penalties',
AGENTIC: 'agentic',
DEVELOPER: 'developer',
TOOLS: 'tools',
IMPORT_EXPORT: 'import-export'
} as const;
export const ROUTES = {
/** Root — start of the app. */
START: '#/',
/** New chat — root with new chat query param. */
NEW_CHAT: `?${NEW_CHAT_PARAM}=true#/`,
/** Chat base — for dynamic chat URLs use RouterService. */
CHAT: '#/chat',
/** MCP servers. */
MCP_SERVERS: '#/mcp-servers',
/** Settings base — for dynamic settings URLs use RouterService. */
SETTINGS: '#/settings'
} as const;

View file

@ -1,163 +0,0 @@
import { ColorMode } from '$lib/enums/ui';
import { Monitor, Moon, Sun } from '@lucide/svelte';
export const SETTING_CONFIG_DEFAULT: Record<string, string | number | boolean | undefined> = {
// Note: in order not to introduce breaking changes, please keep the same data type (number, string, etc) if you want to change the default value.
// Do not use nested objects, keep it single level. Prefix the key if you need to group them.
apiKey: '',
systemMessage: '',
showSystemMessage: true,
theme: ColorMode.SYSTEM,
showThoughtInProgress: true,
disableReasoningParsing: false,
excludeReasoningFromContext: false,
showRawOutputSwitch: false,
keepStatsVisible: false,
showMessageStats: true,
askForTitleConfirmation: false,
titleGenerationUseFirstLine: false,
pasteLongTextToFileLen: 2500,
copyTextAttachmentsAsPlainText: false,
pdfAsImage: false,
disableAutoScroll: false,
renderUserContentAsMarkdown: false,
alwaysShowSidebarOnDesktop: false,
autoShowSidebarOnNewChat: true,
sendOnEnter: true,
autoMicOnEmpty: false,
fullHeightCodeBlocks: false,
showRawModelNames: false,
mcpServers: '[]',
mcpServerUsageStats: '{}', // JSON object: { [serverId]: usageCount }
agenticMaxTurns: 10,
agenticMaxToolPreviewLines: 25,
showToolCallInProgress: false,
alwaysShowAgenticTurns: false,
// sampling params: empty means "use server default"
// the server / preset is the source of truth
// empty values are shown as placeholders from /props in the UI
// and are NOT sent in API requests, letting the server decide
samplers: '',
backend_sampling: false,
temperature: undefined,
dynatemp_range: undefined,
dynatemp_exponent: undefined,
top_k: undefined,
top_p: undefined,
min_p: undefined,
xtc_probability: undefined,
xtc_threshold: undefined,
typ_p: undefined,
repeat_last_n: undefined,
repeat_penalty: undefined,
presence_penalty: undefined,
frequency_penalty: undefined,
dry_multiplier: undefined,
dry_base: undefined,
dry_allowed_length: undefined,
dry_penalty_last_n: undefined,
max_tokens: undefined,
custom: '', // custom json-stringified object
preEncodeConversation: false,
// experimental features
pyInterpreterEnabled: false,
enableContinueGeneration: true
};
export const SETTING_CONFIG_INFO: Record<string, string> = {
apiKey: 'Set the API Key if you are using <code>--api-key</code> option for the server.',
systemMessage: 'The starting message that defines how model should behave.',
showSystemMessage: 'Display the system message at the top of each conversation.',
theme:
'Choose the color theme for the interface. You can choose between System (follows your device settings), Light, or Dark.',
pasteLongTextToFileLen:
'On pasting long text, it will be converted to a file. You can control the file length by setting the value of this parameter. Value 0 means disable.',
copyTextAttachmentsAsPlainText:
'When copying a message with text attachments, combine them into a single plain text string instead of a special format that can be pasted back as attachments.',
samplers:
'The order at which samplers are applied, in simplified way. Default is "top_k;typ_p;top_p;min_p;temperature": top_k->typ_p->top_p->min_p->temperature',
backend_sampling:
'Enable backend-based samplers. When enabled, supported samplers run on the accelerator backend for faster sampling.',
temperature:
'Controls the randomness of the generated text by affecting the probability distribution of the output tokens. Higher = more random, lower = more focused.',
dynatemp_range:
'Addon for the temperature sampler. The added value to the range of dynamic temperature, which adjusts probabilities by entropy of tokens.',
dynatemp_exponent:
'Addon for the temperature sampler. Smoothes out the probability redistribution based on the most probable token.',
top_k: 'Keeps only k top tokens.',
top_p: 'Limits tokens to those that together have a cumulative probability of at least p',
min_p:
'Limits tokens based on the minimum probability for a token to be considered, relative to the probability of the most likely token.',
xtc_probability:
'XTC sampler cuts out top tokens; this parameter controls the chance of cutting tokens at all. 0 disables XTC.',
xtc_threshold:
'XTC sampler cuts out top tokens; this parameter controls the token probability that is required to cut that token.',
typ_p: 'Sorts and limits tokens based on the difference between log-probability and entropy.',
repeat_last_n: 'Last n tokens to consider for penalizing repetition',
repeat_penalty: 'Controls the repetition of token sequences in the generated text',
presence_penalty: 'Limits tokens based on whether they appear in the output or not.',
frequency_penalty: 'Limits tokens based on how often they appear in the output.',
dry_multiplier:
'DRY sampling reduces repetition in generated text even across long contexts. This parameter sets the DRY sampling multiplier.',
dry_base:
'DRY sampling reduces repetition in generated text even across long contexts. This parameter sets the DRY sampling base value.',
dry_allowed_length:
'DRY sampling reduces repetition in generated text even across long contexts. This parameter sets the allowed length for DRY sampling.',
dry_penalty_last_n:
'DRY sampling reduces repetition in generated text even across long contexts. This parameter sets DRY penalty for the last n tokens.',
max_tokens: 'The maximum number of token per output. Use -1 for infinite (no limit).',
custom: 'Custom JSON parameters to send to the API. Must be valid JSON format.',
showThoughtInProgress: 'Expand thought process by default when generating messages.',
disableReasoningParsing:
'Send reasoning_format=none so the server returns thinking tokens inline instead of extracting them into a separate field.',
excludeReasoningFromContext:
'Strip thinking from previous messages before sending. When off, thinking is sent back via the reasoning_content field so the model sees its own chain-of-thought across turns.',
showRawOutputSwitch:
'Show toggle button to display messages as plain text instead of Markdown-formatted content',
keepStatsVisible: 'Keep processing statistics visible after generation finishes.',
showMessageStats:
'Display generation statistics (tokens/second, token count, duration) below each assistant message.',
askForTitleConfirmation:
'Ask for confirmation before automatically changing conversation title when editing the first message.',
titleGenerationUseFirstLine:
'Use only the first non-empty line of the prompt to generate the conversation title.',
pdfAsImage:
'Parse PDF as image instead of text. Automatically falls back to text processing for non-vision models.',
disableAutoScroll:
'Disable automatic scrolling while messages stream so you can control the viewport position manually.',
renderUserContentAsMarkdown: 'Render user messages using markdown formatting in the chat.',
alwaysShowSidebarOnDesktop:
'Always keep the sidebar visible on desktop instead of auto-hiding it.',
autoShowSidebarOnNewChat:
'Automatically show sidebar when starting a new chat. Disable to keep the sidebar hidden until you click on it.',
sendOnEnter:
'Use Enter to send messages and Shift + Enter for new lines. When disabled, use Ctrl/Cmd + Enter.',
autoMicOnEmpty:
'Automatically show microphone button instead of send button when textarea is empty for models with audio modality support.',
fullHeightCodeBlocks:
'Always display code blocks at their full natural height, overriding any height limits.',
showRawModelNames:
'Display full raw model identifiers (e.g. "ggml-org/GLM-4.7-Flash-GGUF:Q8_0") instead of parsed names with badges.',
mcpServers:
'Configure MCP servers as a JSON list. Use the form in the MCP Client settings section to edit.',
mcpServerUsageStats:
'Usage statistics for MCP servers. Tracks how many times tools from each server have been used.',
agenticMaxTurns:
'Maximum number of tool execution cycles before stopping (prevents infinite loops).',
agenticMaxToolPreviewLines:
'Number of lines shown in tool output previews (last N lines). Only these previews and the final LLM response persist after the agentic loop completes.',
showToolCallInProgress:
'Automatically expand tool call details while executing and keep them expanded after completion.',
pyInterpreterEnabled:
'Enable Python interpreter using Pyodide. Allows running Python code in markdown code blocks.',
preEncodeConversation:
'After each response, re-submit the conversation to pre-fill the server KV cache. Makes the next turn faster since the prompt is already encoded while you read the response.',
enableContinueGeneration:
'Enable "Continue" button for assistant messages. Currently works only with non-reasoning models.'
};
export const SETTINGS_COLOR_MODES_CONFIG = [
{ value: ColorMode.SYSTEM, label: 'System', icon: Monitor },
{ value: ColorMode.LIGHT, label: 'Light', icon: Sun },
{ value: ColorMode.DARK, label: 'Dark', icon: Moon }
];

View file

@ -1,33 +0,0 @@
/**
* List of all numeric fields in settings configuration.
* These fields will be converted from strings to numbers during save.
*/
export const NUMERIC_FIELDS = [
'temperature',
'top_k',
'top_p',
'min_p',
'max_tokens',
'pasteLongTextToFileLen',
'dynatemp_range',
'dynatemp_exponent',
'typ_p',
'xtc_probability',
'xtc_threshold',
'repeat_last_n',
'repeat_penalty',
'presence_penalty',
'frequency_penalty',
'dry_multiplier',
'dry_base',
'dry_allowed_length',
'dry_penalty_last_n',
'agenticMaxTurns',
'agenticMaxToolPreviewLines'
] as const;
/**
* Fields that must be positive integers (>= 1).
* These will be clamped to minimum 1 and rounded during save.
*/
export const POSITIVE_INTEGER_FIELDS = ['agenticMaxTurns', 'agenticMaxToolPreviewLines'] as const;

View file

@ -16,6 +16,8 @@ export const SETTINGS_KEYS = {
PDF_AS_IMAGE: 'pdfAsImage',
ASK_FOR_TITLE_CONFIRMATION: 'askForTitleConfirmation',
TITLE_GENERATION_USE_FIRST_LINE: 'titleGenerationUseFirstLine',
TITLE_GENERATION_USE_LLM: 'titleGenerationUseLLM',
TITLE_GENERATION_PROMPT: 'titleGenerationPrompt',
// Display
SHOW_MESSAGE_STATS: 'showMessageStats',
SHOW_THOUGHT_IN_PROGRESS: 'showThoughtInProgress',
@ -26,6 +28,7 @@ export const SETTINGS_KEYS = {
ALWAYS_SHOW_SIDEBAR_ON_DESKTOP: 'alwaysShowSidebarOnDesktop',
FULL_HEIGHT_CODE_BLOCKS: 'fullHeightCodeBlocks',
SHOW_RAW_MODEL_NAMES: 'showRawModelNames',
SHOW_SYSTEM_MESSAGE: 'showSystemMessage',
// Sampling
TEMPERATURE: 'temperature',
DYNATEMP_RANGE: 'dynatemp_range',
@ -49,6 +52,7 @@ export const SETTINGS_KEYS = {
DRY_ALLOWED_LENGTH: 'dry_allowed_length',
DRY_PENALTY_LAST_N: 'dry_penalty_last_n',
// MCP
MCP_SERVERS: 'mcpServers',
AGENTIC_MAX_TURNS: 'agenticMaxTurns',
ALWAYS_SHOW_AGENTIC_TURNS: 'alwaysShowAgenticTurns',
AGENTIC_MAX_TOOL_PREVIEW_LINES: 'agenticMaxToolPreviewLines',
@ -59,5 +63,6 @@ export const SETTINGS_KEYS = {
DISABLE_REASONING_PARSING: 'disableReasoningParsing',
EXCLUDE_REASONING_FROM_CONTEXT: 'excludeReasoningFromContext',
SHOW_RAW_OUTPUT_SWITCH: 'showRawOutputSwitch',
// PY_INTERPRETER_ENABLED: 'pyInterpreterEnabled',
CUSTOM: 'custom'
} as const;

View file

@ -0,0 +1,719 @@
import { ColorMode } from '$lib/enums/ui';
import { SettingsFieldType } from '$lib/enums/settings';
import { SyncableParameterType } from '$lib/enums';
import {
Funnel,
AlertTriangle,
Code,
Monitor,
ListRestart,
Sliders,
PencilRuler,
Database,
Monitor as MonitorIcon,
Sun,
Moon
} from '@lucide/svelte';
import type { Component } from 'svelte';
import type {
SettingsConfigValue,
SyncableParameter,
SettingsEntry,
SettingsSectionTitle,
SettingsSectionEntry,
SettingsSection
} from '$lib/types';
import { SETTINGS_KEYS } from './settings-keys';
import { ROUTES, SETTINGS_SECTION_SLUGS } from './routes';
import { TITLE_GENERATION } from './title-generation';
export const SETTINGS_SECTION_TITLES = {
GENERAL: 'General',
DISPLAY: 'Display',
SAMPLING: 'Sampling',
PENALTIES: 'Penalties',
AGENTIC: 'Agentic',
TOOLS: 'Tools',
IMPORT_EXPORT: 'Import/Export',
DEVELOPER: 'Developer'
} as const;
const STANDALONE_SECTIONS: { title: SettingsSectionTitle; slug: string; icon: Component }[] = [
{ title: SETTINGS_SECTION_TITLES.TOOLS, slug: SETTINGS_SECTION_SLUGS.TOOLS, icon: PencilRuler },
{
title: SETTINGS_SECTION_TITLES.IMPORT_EXPORT,
slug: SETTINGS_SECTION_SLUGS.IMPORT_EXPORT,
icon: Database
}
];
const COLOR_MODE_OPTIONS: Array<{ value: string; label: string; icon: Component }> = [
{ value: ColorMode.SYSTEM, label: 'System', icon: MonitorIcon },
{ value: ColorMode.LIGHT, label: 'Light', icon: Sun },
{ value: ColorMode.DARK, label: 'Dark', icon: Moon }
];
const SETTINGS_REGISTRY: Record<string, SettingsSectionEntry> = {
[SETTINGS_SECTION_SLUGS.GENERAL]: {
title: SETTINGS_SECTION_TITLES.GENERAL,
slug: SETTINGS_SECTION_SLUGS.GENERAL,
icon: Sliders,
settings: [
{
key: SETTINGS_KEYS.THEME,
label: 'Theme',
help: 'Choose the color theme for the interface. You can choose between System (follows your device settings), Light, or Dark.',
defaultValue: ColorMode.SYSTEM,
type: SettingsFieldType.SELECT,
section: SETTINGS_SECTION_SLUGS.GENERAL,
options: COLOR_MODE_OPTIONS,
sync: { serverKey: SETTINGS_KEYS.THEME, paramType: SyncableParameterType.STRING }
},
{
key: SETTINGS_KEYS.API_KEY,
label: 'API Key',
help: 'Set the API Key if you are using <code>--api-key</code> option for the server.',
defaultValue: '',
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.GENERAL
},
{
key: SETTINGS_KEYS.SYSTEM_MESSAGE,
label: 'System Message',
help: 'The starting message that defines how model should behave.',
defaultValue: '',
type: SettingsFieldType.TEXTAREA,
section: SETTINGS_SECTION_SLUGS.GENERAL,
sync: { serverKey: SETTINGS_KEYS.SYSTEM_MESSAGE, paramType: SyncableParameterType.STRING }
},
{
key: SETTINGS_KEYS.PASTE_LONG_TEXT_TO_FILE_LEN,
label: 'Paste long text to file length',
help: 'On pasting long text, it will be converted to a file. You can control the file length by setting the value of this parameter. Value 0 means disable.',
defaultValue: 2500,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.GENERAL,
sync: {
serverKey: SETTINGS_KEYS.PASTE_LONG_TEXT_TO_FILE_LEN,
paramType: SyncableParameterType.NUMBER
}
},
{
key: SETTINGS_KEYS.SEND_ON_ENTER,
label: 'Send message on Enter',
help: 'Use Enter to send messages and Shift + Enter for new lines. When disabled, use Ctrl/Cmd + Enter.',
defaultValue: true,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.GENERAL,
sync: { serverKey: SETTINGS_KEYS.SEND_ON_ENTER, paramType: SyncableParameterType.BOOLEAN }
},
{
key: SETTINGS_KEYS.COPY_TEXT_ATTACHMENTS_AS_PLAIN_TEXT,
label: 'Copy text attachments as plain text',
help: 'When copying a message with text attachments, combine them into a single plain text string instead of a special format that can be pasted back as attachments.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.GENERAL,
sync: {
serverKey: SETTINGS_KEYS.COPY_TEXT_ATTACHMENTS_AS_PLAIN_TEXT,
paramType: SyncableParameterType.BOOLEAN
}
},
{
key: SETTINGS_KEYS.ENABLE_CONTINUE_GENERATION,
label: 'Enable "Continue" button',
help: 'Enable "Continue" button for assistant messages. Currently works only with non-reasoning models.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.GENERAL,
isExperimental: true,
sync: {
serverKey: SETTINGS_KEYS.ENABLE_CONTINUE_GENERATION,
paramType: SyncableParameterType.BOOLEAN
}
},
{
key: SETTINGS_KEYS.PDF_AS_IMAGE,
label: 'Parse PDF as image',
help: 'Parse PDF as image instead of text. Automatically falls back to text processing for non-vision models.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.GENERAL,
sync: { serverKey: SETTINGS_KEYS.PDF_AS_IMAGE, paramType: SyncableParameterType.BOOLEAN }
},
{
key: SETTINGS_KEYS.ASK_FOR_TITLE_CONFIRMATION,
label: 'Ask for confirmation before changing conversation title',
help: 'Ask for confirmation before automatically changing conversation title when editing the first message.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.GENERAL,
sync: {
serverKey: SETTINGS_KEYS.ASK_FOR_TITLE_CONFIRMATION,
paramType: SyncableParameterType.BOOLEAN
}
},
{
key: SETTINGS_KEYS.TITLE_GENERATION_USE_FIRST_LINE,
label: 'Use first non-empty line for conversation title',
help: 'Use only the first non-empty line of the prompt to generate the conversation title.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.GENERAL,
sync: {
serverKey: SETTINGS_KEYS.TITLE_GENERATION_USE_FIRST_LINE,
paramType: SyncableParameterType.BOOLEAN
}
},
{
key: SETTINGS_KEYS.TITLE_GENERATION_USE_LLM,
label: 'Use LLM to generate conversation title',
help: 'Use the LLM to automatically generate conversation titles based on the first message exchange.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.GENERAL,
isExperimental: true
},
{
key: SETTINGS_KEYS.TITLE_GENERATION_PROMPT,
label: 'LLM title generation prompt',
help: 'Optional template for the title generation prompt. Use {{USER}} for the user message and {{ASSISTANT}} for the assistant message.',
defaultValue: TITLE_GENERATION.DEFAULT_PROMPT,
type: SettingsFieldType.TEXTAREA,
section: SETTINGS_SECTION_SLUGS.GENERAL
}
]
},
[SETTINGS_SECTION_SLUGS.DISPLAY]: {
title: SETTINGS_SECTION_TITLES.DISPLAY,
slug: SETTINGS_SECTION_SLUGS.DISPLAY,
icon: Monitor,
settings: [
{
key: SETTINGS_KEYS.SHOW_MESSAGE_STATS,
label: 'Show message generation statistics',
help: 'Display generation statistics (tokens/second, token count, duration) below each assistant message.',
defaultValue: true,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.DISPLAY,
sync: {
serverKey: SETTINGS_KEYS.SHOW_MESSAGE_STATS,
paramType: SyncableParameterType.BOOLEAN
}
},
{
key: SETTINGS_KEYS.SHOW_THOUGHT_IN_PROGRESS,
label: 'Show thought in progress',
help: 'Expand thought process by default when generating messages.',
defaultValue: true,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.DISPLAY,
sync: {
serverKey: SETTINGS_KEYS.SHOW_THOUGHT_IN_PROGRESS,
paramType: SyncableParameterType.BOOLEAN
}
},
{
key: SETTINGS_KEYS.SHOW_TOOL_CALL_IN_PROGRESS,
label: 'Show tool call in progress',
help: 'Automatically expand tool call details while executing and keep them expanded after completion.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.DISPLAY,
sync: {
serverKey: SETTINGS_KEYS.SHOW_TOOL_CALL_IN_PROGRESS,
paramType: SyncableParameterType.BOOLEAN
}
},
{
key: SETTINGS_KEYS.KEEP_STATS_VISIBLE,
label: 'Keep stats visible after generation',
help: 'Keep processing statistics visible after generation finishes.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.DISPLAY,
sync: {
serverKey: SETTINGS_KEYS.KEEP_STATS_VISIBLE,
paramType: SyncableParameterType.BOOLEAN
}
},
{
key: SETTINGS_KEYS.AUTO_MIC_ON_EMPTY,
label: 'Show microphone on empty input',
help: 'Automatically show microphone button instead of send button when textarea is empty for models with audio modality support.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.DISPLAY,
isExperimental: true,
sync: {
serverKey: SETTINGS_KEYS.AUTO_MIC_ON_EMPTY,
paramType: SyncableParameterType.BOOLEAN
}
},
{
key: SETTINGS_KEYS.RENDER_USER_CONTENT_AS_MARKDOWN,
label: 'Render user content as Markdown',
help: 'Render user messages using markdown formatting in the chat.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.DISPLAY,
sync: {
serverKey: SETTINGS_KEYS.RENDER_USER_CONTENT_AS_MARKDOWN,
paramType: SyncableParameterType.BOOLEAN
}
},
{
key: SETTINGS_KEYS.FULL_HEIGHT_CODE_BLOCKS,
label: 'Use full height code blocks',
help: 'Always display code blocks at their full natural height, overriding any height limits.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.DISPLAY,
sync: {
serverKey: SETTINGS_KEYS.FULL_HEIGHT_CODE_BLOCKS,
paramType: SyncableParameterType.BOOLEAN
}
},
{
key: SETTINGS_KEYS.DISABLE_AUTO_SCROLL,
label: 'Disable automatic scroll',
help: 'Disable automatic scrolling while messages stream so you can control the viewport position manually.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.DISPLAY,
sync: {
serverKey: SETTINGS_KEYS.DISABLE_AUTO_SCROLL,
paramType: SyncableParameterType.BOOLEAN
}
},
{
key: SETTINGS_KEYS.ALWAYS_SHOW_SIDEBAR_ON_DESKTOP,
label: 'Always show sidebar on desktop',
help: 'Always keep the sidebar visible on desktop instead of auto-hiding it.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.DISPLAY,
sync: {
serverKey: SETTINGS_KEYS.ALWAYS_SHOW_SIDEBAR_ON_DESKTOP,
paramType: SyncableParameterType.BOOLEAN
}
},
{
key: SETTINGS_KEYS.SHOW_RAW_MODEL_NAMES,
label: 'Show raw model names',
help: 'Display full raw model identifiers (e.g. "ggml-org/GLM-4.7-Flash-GGUF:Q8_0") instead of parsed names with badges.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.DISPLAY,
sync: {
serverKey: SETTINGS_KEYS.SHOW_RAW_MODEL_NAMES,
paramType: SyncableParameterType.BOOLEAN
}
},
{
key: SETTINGS_KEYS.ALWAYS_SHOW_AGENTIC_TURNS,
label: 'Always show agentic turns in conversation',
help: 'Always expand and display agentic loop turns in conversation messages.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.DISPLAY,
sync: {
serverKey: SETTINGS_KEYS.ALWAYS_SHOW_AGENTIC_TURNS,
paramType: SyncableParameterType.BOOLEAN
}
}
]
},
[SETTINGS_SECTION_SLUGS.SAMPLING]: {
title: SETTINGS_SECTION_TITLES.SAMPLING,
slug: SETTINGS_SECTION_SLUGS.SAMPLING,
icon: Funnel,
settings: [
{
key: SETTINGS_KEYS.TEMPERATURE,
label: 'Temperature',
help: 'Controls the randomness of the generated text by affecting the probability distribution of the output tokens. Higher = more random, lower = more focused.',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.SAMPLING,
sync: { serverKey: SETTINGS_KEYS.TEMPERATURE, paramType: SyncableParameterType.NUMBER }
},
{
key: SETTINGS_KEYS.DYNATEMP_RANGE,
label: 'Dynamic temperature range',
help: 'Addon for the temperature sampler. The added value to the range of dynamic temperature, which adjusts probabilities by entropy of tokens.',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.SAMPLING,
sync: { serverKey: SETTINGS_KEYS.DYNATEMP_RANGE, paramType: SyncableParameterType.NUMBER }
},
{
key: SETTINGS_KEYS.DYNATEMP_EXPONENT,
label: 'Dynamic temperature exponent',
help: 'Addon for the temperature sampler. Smoothes out the probability redistribution based on the most probable token.',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.SAMPLING,
sync: {
serverKey: SETTINGS_KEYS.DYNATEMP_EXPONENT,
paramType: SyncableParameterType.NUMBER
}
},
{
key: SETTINGS_KEYS.TOP_K,
label: 'Top K',
help: 'Keeps only k top tokens.',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.SAMPLING,
sync: { serverKey: SETTINGS_KEYS.TOP_K, paramType: SyncableParameterType.NUMBER }
},
{
key: SETTINGS_KEYS.TOP_P,
label: 'Top P',
help: 'Limits tokens to those that together have a cumulative probability of at least p',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.SAMPLING,
sync: { serverKey: SETTINGS_KEYS.TOP_P, paramType: SyncableParameterType.NUMBER }
},
{
key: SETTINGS_KEYS.MIN_P,
label: 'Min P',
help: 'Limits tokens based on the minimum probability for a token to be considered, relative to the probability of the most likely token.',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.SAMPLING,
sync: { serverKey: SETTINGS_KEYS.MIN_P, paramType: SyncableParameterType.NUMBER }
},
{
key: SETTINGS_KEYS.XTC_PROBABILITY,
label: 'XTC probability',
help: 'XTC sampler cuts out top tokens; this parameter controls the chance of cutting tokens at all. 0 disables XTC.',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.SAMPLING,
sync: { serverKey: SETTINGS_KEYS.XTC_PROBABILITY, paramType: SyncableParameterType.NUMBER }
},
{
key: SETTINGS_KEYS.XTC_THRESHOLD,
label: 'XTC threshold',
help: 'XTC sampler cuts out top tokens; this parameter controls the token probability that is required to cut that token.',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.SAMPLING,
sync: { serverKey: SETTINGS_KEYS.XTC_THRESHOLD, paramType: SyncableParameterType.NUMBER }
},
{
key: SETTINGS_KEYS.TYP_P,
label: 'Typical P',
help: 'Sorts and limits tokens based on the difference between log-probability and entropy.',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.SAMPLING,
sync: { serverKey: SETTINGS_KEYS.TYP_P, paramType: SyncableParameterType.NUMBER }
},
{
key: SETTINGS_KEYS.MAX_TOKENS,
label: 'Max tokens',
help: 'The maximum number of token per output. Use -1 for infinite (no limit).',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.SAMPLING,
sync: { serverKey: SETTINGS_KEYS.MAX_TOKENS, paramType: SyncableParameterType.NUMBER }
},
{
key: SETTINGS_KEYS.SAMPLERS,
label: 'Samplers',
help: 'The order at which samplers are applied, in simplified way. Default is "top_k;typ_p;top_p;min_p;temperature": top_k->typ_p->top_p->min_p->temperature',
defaultValue: '',
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.SAMPLING,
sync: { serverKey: SETTINGS_KEYS.SAMPLERS, paramType: SyncableParameterType.STRING }
},
{
key: SETTINGS_KEYS.BACKEND_SAMPLING,
label: 'Backend sampling',
help: 'Enable backend-based samplers. When enabled, supported samplers run on the accelerator backend for faster sampling.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.SAMPLING,
sync: {
serverKey: SETTINGS_KEYS.BACKEND_SAMPLING,
paramType: SyncableParameterType.BOOLEAN
}
}
]
},
[SETTINGS_SECTION_SLUGS.PENALTIES]: {
title: SETTINGS_SECTION_TITLES.PENALTIES,
slug: SETTINGS_SECTION_SLUGS.PENALTIES,
icon: AlertTriangle,
settings: [
{
key: SETTINGS_KEYS.REPEAT_LAST_N,
label: 'Repeat last N',
help: 'Last n tokens to consider for penalizing repetition',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.PENALTIES,
sync: { serverKey: SETTINGS_KEYS.REPEAT_LAST_N, paramType: SyncableParameterType.NUMBER }
},
{
key: SETTINGS_KEYS.REPEAT_PENALTY,
label: 'Repeat penalty',
help: 'Controls the repetition of token sequences in the generated text',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.PENALTIES,
sync: { serverKey: SETTINGS_KEYS.REPEAT_PENALTY, paramType: SyncableParameterType.NUMBER }
},
{
key: SETTINGS_KEYS.PRESENCE_PENALTY,
label: 'Presence penalty',
help: 'Limits tokens based on whether they appear in the output or not.',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.PENALTIES,
sync: { serverKey: SETTINGS_KEYS.PRESENCE_PENALTY, paramType: SyncableParameterType.NUMBER }
},
{
key: SETTINGS_KEYS.FREQUENCY_PENALTY,
label: 'Frequency penalty',
help: 'Limits tokens based on how often they appear in the output.',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.PENALTIES,
sync: {
serverKey: SETTINGS_KEYS.FREQUENCY_PENALTY,
paramType: SyncableParameterType.NUMBER
}
},
{
key: SETTINGS_KEYS.DRY_MULTIPLIER,
label: 'DRY multiplier',
help: 'DRY sampling reduces repetition in generated text even across long contexts. This parameter sets the DRY sampling multiplier.',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.PENALTIES,
sync: { serverKey: SETTINGS_KEYS.DRY_MULTIPLIER, paramType: SyncableParameterType.NUMBER }
},
{
key: SETTINGS_KEYS.DRY_BASE,
label: 'DRY base',
help: 'DRY sampling reduces repetition in generated text even across long contexts. This parameter sets the DRY sampling base value.',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.PENALTIES,
sync: { serverKey: SETTINGS_KEYS.DRY_BASE, paramType: SyncableParameterType.NUMBER }
},
{
key: SETTINGS_KEYS.DRY_ALLOWED_LENGTH,
label: 'DRY allowed length',
help: 'DRY sampling reduces repetition in generated text even across long contexts. This parameter sets the allowed length for DRY sampling.',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.PENALTIES,
sync: {
serverKey: SETTINGS_KEYS.DRY_ALLOWED_LENGTH,
paramType: SyncableParameterType.NUMBER
}
},
{
key: SETTINGS_KEYS.DRY_PENALTY_LAST_N,
label: 'DRY penalty last N',
help: 'DRY sampling reduces repetition in generated text even across long contexts. This parameter sets DRY penalty for the last n tokens.',
defaultValue: undefined,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.PENALTIES,
sync: {
serverKey: SETTINGS_KEYS.DRY_PENALTY_LAST_N,
paramType: SyncableParameterType.NUMBER
}
}
]
},
[SETTINGS_SECTION_SLUGS.AGENTIC]: {
title: SETTINGS_SECTION_TITLES.AGENTIC,
slug: SETTINGS_SECTION_SLUGS.AGENTIC,
icon: ListRestart,
settings: [
{
key: SETTINGS_KEYS.AGENTIC_MAX_TURNS,
label: 'Agentic turns',
help: 'Maximum number of tool execution cycles before stopping (prevents infinite loops).',
defaultValue: 10,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.AGENTIC,
isPositiveInteger: true,
sync: {
serverKey: SETTINGS_KEYS.AGENTIC_MAX_TURNS,
paramType: SyncableParameterType.NUMBER
}
},
{
key: SETTINGS_KEYS.AGENTIC_MAX_TOOL_PREVIEW_LINES,
label: 'Max lines per tool preview',
help: 'Number of lines shown in tool output previews (last N lines). Only these previews and the final LLM response persist after the agentic loop completes.',
defaultValue: 25,
type: SettingsFieldType.INPUT,
section: SETTINGS_SECTION_SLUGS.AGENTIC,
isPositiveInteger: true,
sync: {
serverKey: SETTINGS_KEYS.AGENTIC_MAX_TOOL_PREVIEW_LINES,
paramType: SyncableParameterType.NUMBER
}
}
]
},
[SETTINGS_SECTION_SLUGS.DEVELOPER]: {
title: SETTINGS_SECTION_TITLES.DEVELOPER,
slug: SETTINGS_SECTION_SLUGS.DEVELOPER,
icon: Code,
settings: [
{
key: SETTINGS_KEYS.PRE_ENCODE_CONVERSATION,
label: 'Pre-fill KV cache after response',
help: 'After each response, re-submit the conversation to pre-fill the server KV cache. Makes the next turn faster since the prompt is already encoded while you read the response.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.DEVELOPER
},
{
key: SETTINGS_KEYS.DISABLE_REASONING_PARSING,
label: 'Disable reasoning content parsing',
help: 'Send reasoning_format=none so the server returns thinking tokens inline instead of extracting them into a separate field.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.DEVELOPER
},
{
key: SETTINGS_KEYS.EXCLUDE_REASONING_FROM_CONTEXT,
label: 'Exclude reasoning from context',
help: 'Strip thinking from previous messages before sending. When off, thinking is sent back via the reasoning_content field so the model sees its own chain-of-thought across turns.',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.DEVELOPER,
sync: {
serverKey: SETTINGS_KEYS.EXCLUDE_REASONING_FROM_CONTEXT,
paramType: SyncableParameterType.BOOLEAN
}
},
{
key: SETTINGS_KEYS.SHOW_RAW_OUTPUT_SWITCH,
label: 'Enable raw output toggle',
help: 'Show toggle button to display messages as plain text instead of Markdown-formatted content',
defaultValue: false,
type: SettingsFieldType.CHECKBOX,
section: SETTINGS_SECTION_SLUGS.DEVELOPER,
sync: {
serverKey: SETTINGS_KEYS.SHOW_RAW_OUTPUT_SWITCH,
paramType: SyncableParameterType.BOOLEAN
}
},
{
key: SETTINGS_KEYS.CUSTOM,
label: 'Custom JSON',
help: 'Custom JSON parameters to send to the API. Must be valid JSON format.',
defaultValue: '',
type: SettingsFieldType.TEXTAREA,
section: SETTINGS_SECTION_SLUGS.DEVELOPER
}
]
}
} as const;
const NON_UI_SETTINGS: SettingsEntry[] = [
{
key: SETTINGS_KEYS.SHOW_SYSTEM_MESSAGE,
label: 'Show system message',
help: 'Display the system message at the top of each conversation.',
defaultValue: true,
type: SettingsFieldType.CHECKBOX,
sync: { serverKey: SETTINGS_KEYS.SHOW_SYSTEM_MESSAGE, paramType: SyncableParameterType.BOOLEAN }
},
{
key: SETTINGS_KEYS.MCP_SERVERS,
label: 'MCP servers',
help: 'Configure MCP servers as a JSON list. Use the form in the MCP Client settings section to edit.',
defaultValue: '[]',
type: SettingsFieldType.INPUT,
sync: { serverKey: SETTINGS_KEYS.MCP_SERVERS, paramType: SyncableParameterType.STRING }
}
// {
// key: SETTINGS_KEYS.PY_INTERPRETER_ENABLED,
// label: 'Python interpreter enabled',
// help: 'Enable Python interpreter using Pyodide. Allows running Python code in markdown code blocks.',
// defaultValue: false,
// type: SettingsFieldType.CHECKBOX,
// isExperimental: true,
// sync: { serverKey: SETTINGS_KEYS.PY_INTERPRETER_ENABLED, paramType: SyncableParameterType.BOOLEAN }
// }
];
function getAllSettings(): SettingsEntry[] {
const result: SettingsEntry[] = [];
for (const section of Object.values(SETTINGS_REGISTRY)) {
result.push(...section.settings);
}
result.push(...NON_UI_SETTINGS);
return result;
}
/** Flat config object stored in localStorage. */
export const SETTING_CONFIG_DEFAULT: Record<string, SettingsConfigValue> = Object.fromEntries(
getAllSettings().map((s) => [s.key, s.defaultValue])
) as Record<string, SettingsConfigValue>;
/** Help text for every setting (including non-UI). */
export const SETTING_CONFIG_INFO: Record<string, string> = Object.fromEntries(
getAllSettings().map((s) => [s.key, s.help])
) as Record<string, string>;
/** Theme select options. */
export const SETTINGS_COLOR_MODES_CONFIG = COLOR_MODE_OPTIONS;
export type { SettingsSectionTitle } from '$lib/types';
export type { SettingsSection } from '$lib/types';
/** Sidebar sections + field configs (as consumed by UI). */
export const SETTINGS_CHAT_SECTIONS: SettingsSection[] = [
...Object.values(SETTINGS_REGISTRY).map((section) => ({
title: section.title,
slug: section.slug,
icon: section.icon,
fields: section.settings.map((s) => ({
key: s.key,
label: s.label,
type: s.type,
isExperimental: s.isExperimental,
help: s.help,
options: s.options
}))
})),
...STANDALONE_SECTIONS
];
/** INPUT-type settings whose value is a number. */
export const NUMERIC_FIELDS = getAllSettings()
.filter((s) => s.type === SettingsFieldType.INPUT && typeof s.defaultValue !== 'string')
.map((s) => s.key) as readonly string[];
/** Numeric fields clamped to ≥ 1 and rounded. */
export const POSITIVE_INTEGER_FIELDS = getAllSettings()
.filter((s) => s.isPositiveInteger)
.map((s) => s.key) as readonly string[];
/** Derived for the parameter sync service. */
export const SYNCABLE_PARAMETERS: SyncableParameter[] = getAllSettings()
.filter((s) => s.sync !== undefined)
.map((s) => ({
key: s.key,
serverKey: s.sync!.serverKey,
type: s.sync!.paramType,
canSync: true
}));
export const SETTINGS_FALLBACK_EXIT_ROUTE = ROUTES.START;
export { SETTINGS_KEYS } from './settings-keys';

View file

@ -1,337 +0,0 @@
/**
* Settings section titles constants for ChatSettings component.
*
* These titles define the navigation sections in the settings dialog.
* Used for both sidebar navigation and mobile horizontal scroll menu.
*/
export const SETTINGS_SECTION_TITLES = {
GENERAL: 'General',
DISPLAY: 'Display',
SAMPLING: 'Sampling',
PENALTIES: 'Penalties',
AGENTIC: 'Agentic',
TOOLS: 'Tools',
MCP: 'MCP',
IMPORT_EXPORT: 'Import/Export',
DEVELOPER: 'Developer'
} as const;
/** Type for settings section titles */
export type SettingsSectionTitle =
(typeof SETTINGS_SECTION_TITLES)[keyof typeof SETTINGS_SECTION_TITLES];
import {
Funnel,
AlertTriangle,
Code,
Monitor,
ListRestart,
Sliders,
PencilRuler,
Database
} from '@lucide/svelte';
import { SettingsFieldType } from '$lib/enums/settings';
import { SETTINGS_COLOR_MODES_CONFIG } from '$lib/constants/settings-config';
import { SETTINGS_KEYS } from '$lib/constants/settings-keys';
import type { Component } from 'svelte';
export interface SettingsSection {
fields?: SettingsFieldConfig[];
icon: Component;
slug: string;
title: SettingsSectionTitle;
}
export const SETTINGS_CHAT_SECTIONS: SettingsSection[] = [
{
title: SETTINGS_SECTION_TITLES.GENERAL,
slug: 'general',
icon: Sliders,
fields: [
{
key: SETTINGS_KEYS.THEME,
label: 'Theme',
type: SettingsFieldType.SELECT,
options: SETTINGS_COLOR_MODES_CONFIG
},
{ key: SETTINGS_KEYS.API_KEY, label: 'API Key', type: SettingsFieldType.INPUT },
{
key: SETTINGS_KEYS.SYSTEM_MESSAGE,
label: 'System Message',
type: SettingsFieldType.TEXTAREA
},
{
key: SETTINGS_KEYS.PASTE_LONG_TEXT_TO_FILE_LEN,
label: 'Paste long text to file length',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.SEND_ON_ENTER,
label: 'Send message on Enter',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.COPY_TEXT_ATTACHMENTS_AS_PLAIN_TEXT,
label: 'Copy text attachments as plain text',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.ENABLE_CONTINUE_GENERATION,
label: 'Enable "Continue" button',
type: SettingsFieldType.CHECKBOX,
isExperimental: true
},
{
key: SETTINGS_KEYS.PDF_AS_IMAGE,
label: 'Parse PDF as image',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.ASK_FOR_TITLE_CONFIRMATION,
label: 'Ask for confirmation before changing conversation title',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.TITLE_GENERATION_USE_FIRST_LINE,
label: 'Use first non-empty line for conversation title',
type: SettingsFieldType.CHECKBOX
}
]
},
{
title: SETTINGS_SECTION_TITLES.DISPLAY,
slug: 'display',
icon: Monitor,
fields: [
{
key: SETTINGS_KEYS.SHOW_MESSAGE_STATS,
label: 'Show message generation statistics',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.SHOW_THOUGHT_IN_PROGRESS,
label: 'Show thought in progress',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.SHOW_TOOL_CALL_IN_PROGRESS,
label: 'Show tool call in progress',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.KEEP_STATS_VISIBLE,
label: 'Keep stats visible after generation',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.AUTO_MIC_ON_EMPTY,
label: 'Show microphone on empty input',
type: SettingsFieldType.CHECKBOX,
isExperimental: true
},
{
key: SETTINGS_KEYS.RENDER_USER_CONTENT_AS_MARKDOWN,
label: 'Render user content as Markdown',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.FULL_HEIGHT_CODE_BLOCKS,
label: 'Use full height code blocks',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.DISABLE_AUTO_SCROLL,
label: 'Disable automatic scroll',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.ALWAYS_SHOW_SIDEBAR_ON_DESKTOP,
label: 'Always show sidebar on desktop',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.SHOW_RAW_MODEL_NAMES,
label: 'Show raw model names',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.ALWAYS_SHOW_AGENTIC_TURNS,
label: 'Always show agentic turns in conversation',
type: SettingsFieldType.CHECKBOX
}
]
},
{
title: SETTINGS_SECTION_TITLES.SAMPLING,
slug: 'sampling',
icon: Funnel,
fields: [
{
key: SETTINGS_KEYS.TEMPERATURE,
label: 'Temperature',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.DYNATEMP_RANGE,
label: 'Dynamic temperature range',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.DYNATEMP_EXPONENT,
label: 'Dynamic temperature exponent',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.TOP_K,
label: 'Top K',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.TOP_P,
label: 'Top P',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.MIN_P,
label: 'Min P',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.XTC_PROBABILITY,
label: 'XTC probability',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.XTC_THRESHOLD,
label: 'XTC threshold',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.TYP_P,
label: 'Typical P',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.MAX_TOKENS,
label: 'Max tokens',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.SAMPLERS,
label: 'Samplers',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.BACKEND_SAMPLING,
label: 'Backend sampling',
type: SettingsFieldType.CHECKBOX
}
]
},
{
title: SETTINGS_SECTION_TITLES.PENALTIES,
slug: 'penalties',
icon: AlertTriangle,
fields: [
{
key: SETTINGS_KEYS.REPEAT_LAST_N,
label: 'Repeat last N',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.REPEAT_PENALTY,
label: 'Repeat penalty',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.PRESENCE_PENALTY,
label: 'Presence penalty',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.FREQUENCY_PENALTY,
label: 'Frequency penalty',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.DRY_MULTIPLIER,
label: 'DRY multiplier',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.DRY_BASE,
label: 'DRY base',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.DRY_ALLOWED_LENGTH,
label: 'DRY allowed length',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.DRY_PENALTY_LAST_N,
label: 'DRY penalty last N',
type: SettingsFieldType.INPUT
}
]
},
{
title: SETTINGS_SECTION_TITLES.AGENTIC,
slug: 'agentic',
icon: ListRestart,
fields: [
{
key: SETTINGS_KEYS.AGENTIC_MAX_TURNS,
label: 'Agentic turns',
type: SettingsFieldType.INPUT
},
{
key: SETTINGS_KEYS.AGENTIC_MAX_TOOL_PREVIEW_LINES,
label: 'Max lines per tool preview',
type: SettingsFieldType.INPUT
}
]
},
{
title: SETTINGS_SECTION_TITLES.TOOLS,
slug: 'tools',
icon: PencilRuler
},
{
title: SETTINGS_SECTION_TITLES.IMPORT_EXPORT,
slug: 'import-export',
icon: Database
},
{
title: SETTINGS_SECTION_TITLES.DEVELOPER,
slug: 'developer',
icon: Code,
fields: [
{
key: SETTINGS_KEYS.PRE_ENCODE_CONVERSATION,
label: 'Pre-fill KV cache after response',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.DISABLE_REASONING_PARSING,
label: 'Disable reasoning content parsing',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.EXCLUDE_REASONING_FROM_CONTEXT,
label: 'Exclude reasoning from context',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.SHOW_RAW_OUTPUT_SWITCH,
label: 'Enable raw output toggle',
type: SettingsFieldType.CHECKBOX
},
{
key: SETTINGS_KEYS.CUSTOM,
label: 'Custom JSON',
type: SettingsFieldType.TEXTAREA
}
]
}
];

View file

@ -0,0 +1,9 @@
/* Title generation constants */
export const TITLE_GENERATION = {
MIN_LENGTH: 3,
FALLBACK: 'New Chat',
DEFAULT_PROMPT:
'Based on the following interaction, generate a short, concise title (maximum 6-8 words) that captures the main topic. Return ONLY the title text, nothing else. Do not use quotes.\n\nUser: {{USER}}\n\nAssistant: {{ASSISTANT}}\n\nTitle:',
PREFIX_PATTERN: /^(Title:|Subject:|Topic:)\s*/i,
QUOTE_PATTERN: /^["]|["]$/g
} as const;

View file

@ -1,6 +1,7 @@
import { Settings, Search, SquarePen } from '@lucide/svelte';
import McpLogo from '$lib/components/app/mcp/McpLogo.svelte';
import type { Component } from 'svelte';
import { ROUTES } from './routes';
export const FORK_TREE_DEPTH_PADDING = 8;
export const SYSTEM_MESSAGE_PLACEHOLDER = 'System message';
@ -19,18 +20,18 @@ export interface DesktopIconStripItem {
}
export const SIDEBAR_ACTIONS_ITEMS: DesktopIconStripItem[] = [
{ icon: SquarePen, tooltip: 'New chat', route: '?new_chat=true#/', keys: ['shift', 'cmd', 'o'] },
{ icon: SquarePen, tooltip: 'New chat', route: ROUTES.NEW_CHAT, keys: ['shift', 'cmd', 'o'] },
{ icon: Search, tooltip: 'Search', keys: ['cmd', 'k'] },
{
icon: McpLogo,
tooltip: 'MCP Servers',
route: '#/settings/mcp',
activeRouteId: '/settings/mcp'
route: ROUTES.MCP_SERVERS,
activeRouteId: '/mcp-servers'
},
{
icon: Settings,
tooltip: 'Settings',
route: '#/settings/chat/general',
activeRoutePrefix: '/settings/chat'
route: ROUTES.SETTINGS,
activeRoutePrefix: '/settings'
}
];

View file

@ -1,5 +1,6 @@
import { goto } from '$app/navigation';
import { KeyboardKey } from '$lib/enums';
import { ROUTES } from '$lib/constants/routes';
interface KeyboardShortcutsCallbacks {
activateSearchMode?: () => void;
@ -27,7 +28,7 @@ export function useKeyboardShortcuts(callbacks: KeyboardShortcutsCallbacks) {
) {
event.preventDefault();
goto('?new_chat=true#/');
goto(ROUTES.NEW_CHAT);
}
if (event.shiftKey && isCmdOrCtrl && event.key === KeyboardKey.E_UPPER) {

View file

@ -1,6 +1,7 @@
import { page } from '$app/state';
import { beforeNavigate } from '$app/navigation';
import { settingsReferrer } from '$lib/stores/settings-referrer.svelte';
import { ROUTES } from '$lib/constants/routes';
export interface ChatSettings {
reset: () => void;
@ -15,11 +16,8 @@ export function useSettingsNavigation() {
const isSettingsRoute = $derived(!!page.route.id?.startsWith('/settings'));
beforeNavigate(({ to, from }) => {
if (
to?.route?.id?.startsWith('/settings/chat') &&
!from?.route?.id?.startsWith('/settings/chat')
) {
settingsReferrer.url = window.location.hash || '#/';
if (to?.route?.id?.startsWith('/settings') && !from?.route?.id?.startsWith('/settings')) {
settingsReferrer.url = window.location.hash || ROUTES.START;
}
});

View file

@ -14,11 +14,59 @@ import {
ReasoningFormat,
UrlProtocol
} from '$lib/enums';
import type { ApiChatMessageContentPart, ApiChatCompletionToolCall } from '$lib/types/api';
import type {
ApiChatMessageContentPart,
ApiChatMessageData,
ApiChatCompletionToolCall
} from '$lib/types/api';
import type { DatabaseMessageExtraMcpPrompt, DatabaseMessageExtraMcpResource } from '$lib/types';
import { modelsStore } from '$lib/stores/models.svelte';
export class ChatService {
/**
*
*
* Title Generation
*
*
*/
/**
* Sends a streaming chat completion request for generating a chat title.
* Delegates to `sendMessage` for fetch, SSE parsing, and error handling.
*
* @param message - The single message to send (a user message containing the title generation prompt)
* @param model - Optional model name to use (required in ROUTER mode)
* @param signal - Optional AbortSignal to cancel the request
* @returns {Promise<string>} The aggregated title text, or empty string if request failed
* @static
*/
static async generateTitle(
message: ApiChatMessageData,
model?: string | null,
signal?: AbortSignal
): Promise<string> {
let titleResponse = '';
try {
await ChatService.sendMessage(
[message],
{
model: model || undefined,
stream: true,
custom: { chat_template_kwargs: { enable_thinking: false } },
onChunk: (chunk: string) => {
titleResponse += chunk;
}
},
undefined,
signal
);
} catch {
return '';
}
return titleResponse;
}
/**
*
*
@ -122,7 +170,11 @@ export class ChatService {
return true;
});
// If only text remains and it's a single part, simplify to string
if (msg.content.length === 1 && msg.content[0].type === ContentPartType.TEXT) {
if (
msg.content.length === 1 &&
msg.content[0].type === ContentPartType.TEXT &&
typeof msg.content[0].text === 'string'
) {
msg.content = msg.content[0].text;
}
}
@ -461,7 +513,7 @@ export class ChatService {
const serializedToolCalls = JSON.stringify(aggregatedToolCalls);
if (import.meta.env.DEV) {
if (import.meta.env.DEV && import.meta.env.VITE_DEBUG) {
console.log('[ChatService] Aggregated tool calls:', serializedToolCalls);
}

View file

@ -260,3 +260,26 @@ export { ParameterSyncService } from './parameter-sync.service';
* @see MCP Protocol Specification: https://modelcontextprotocol.io/specification/2025-06-18
*/
export { MCPService } from './mcp.service';
/**
* **RouterService** Dynamic route URL construction utility
*
* Stateless utility for building dynamic route URLs from ROUTES base paths.
* Static routes (START, NEW_CHAT, MCP_SERVERS) live in ROUTES constants;
* dynamic routes (CHAT, SETTINGS) are constructed here by appending parameters.
*
* **Architecture & Relationships:**
* - **RouterService** (this class): Stateless URL construction
* - Builds dynamic route URLs from ROUTES base paths
* - No side effects receives route parameters, returns route strings
*
* - **ROUTES constant** (constants/routes.ts): Static route base paths
* - **All components/stores**: Call RouterService for dynamic route URLs
*
* **Key Responsibilities:**
* - Build chat URLs for specific conversations: `RouterService.chat(id)` `#/chat/:id`
* - Build settings URLs for sections: `RouterService.settings(section)` `#/settings/:section`
*
* @see ROUTES in constants/routes.ts static route base paths
*/
export { RouterService } from './router.service';

View file

@ -1,253 +1,8 @@
import { normalizeFloatingPoint } from '$lib/utils';
import type { SyncableParameter, ParameterRecord, ParameterInfo, ParameterValue } from '$lib/types';
import { SETTINGS_KEYS, SYNCABLE_PARAMETERS } from '$lib/constants';
import type { ParameterRecord, ParameterInfo, ParameterValue } from '$lib/types';
import { SyncableParameterType, ParameterSource } from '$lib/enums';
/**
* Mapping of webui setting keys to server parameter keys.
* Only parameters listed here can be synced from the server `/props` endpoint.
* Each entry defines the webui key, corresponding server key, value type,
* and whether sync is enabled.
*/
export const SYNCABLE_PARAMETERS: SyncableParameter[] = [
{
key: 'temperature',
serverKey: 'temperature',
type: SyncableParameterType.NUMBER,
canSync: true
},
{ key: 'top_k', serverKey: 'top_k', type: SyncableParameterType.NUMBER, canSync: true },
{ key: 'top_p', serverKey: 'top_p', type: SyncableParameterType.NUMBER, canSync: true },
{ key: 'min_p', serverKey: 'min_p', type: SyncableParameterType.NUMBER, canSync: true },
{
key: 'dynatemp_range',
serverKey: 'dynatemp_range',
type: SyncableParameterType.NUMBER,
canSync: true
},
{
key: 'dynatemp_exponent',
serverKey: 'dynatemp_exponent',
type: SyncableParameterType.NUMBER,
canSync: true
},
{
key: 'xtc_probability',
serverKey: 'xtc_probability',
type: SyncableParameterType.NUMBER,
canSync: true
},
{
key: 'xtc_threshold',
serverKey: 'xtc_threshold',
type: SyncableParameterType.NUMBER,
canSync: true
},
{ key: 'typ_p', serverKey: 'typ_p', type: SyncableParameterType.NUMBER, canSync: true },
{
key: 'repeat_last_n',
serverKey: 'repeat_last_n',
type: SyncableParameterType.NUMBER,
canSync: true
},
{
key: 'repeat_penalty',
serverKey: 'repeat_penalty',
type: SyncableParameterType.NUMBER,
canSync: true
},
{
key: 'presence_penalty',
serverKey: 'presence_penalty',
type: SyncableParameterType.NUMBER,
canSync: true
},
{
key: 'frequency_penalty',
serverKey: 'frequency_penalty',
type: SyncableParameterType.NUMBER,
canSync: true
},
{
key: 'dry_multiplier',
serverKey: 'dry_multiplier',
type: SyncableParameterType.NUMBER,
canSync: true
},
{ key: 'dry_base', serverKey: 'dry_base', type: SyncableParameterType.NUMBER, canSync: true },
{
key: 'dry_allowed_length',
serverKey: 'dry_allowed_length',
type: SyncableParameterType.NUMBER,
canSync: true
},
{
key: 'dry_penalty_last_n',
serverKey: 'dry_penalty_last_n',
type: SyncableParameterType.NUMBER,
canSync: true
},
{ key: 'max_tokens', serverKey: 'max_tokens', type: SyncableParameterType.NUMBER, canSync: true },
{ key: 'samplers', serverKey: 'samplers', type: SyncableParameterType.STRING, canSync: true },
{
key: 'backend_sampling',
serverKey: 'backend_sampling',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'pasteLongTextToFileLen',
serverKey: 'pasteLongTextToFileLen',
type: SyncableParameterType.NUMBER,
canSync: true
},
{
key: 'pdfAsImage',
serverKey: 'pdfAsImage',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'showThoughtInProgress',
serverKey: 'showThoughtInProgress',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'keepStatsVisible',
serverKey: 'keepStatsVisible',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'showMessageStats',
serverKey: 'showMessageStats',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'askForTitleConfirmation',
serverKey: 'askForTitleConfirmation',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'titleGenerationUseFirstLine',
serverKey: 'titleGenerationUseFirstLine',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'disableAutoScroll',
serverKey: 'disableAutoScroll',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'renderUserContentAsMarkdown',
serverKey: 'renderUserContentAsMarkdown',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'autoMicOnEmpty',
serverKey: 'autoMicOnEmpty',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'pyInterpreterEnabled',
serverKey: 'pyInterpreterEnabled',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'enableContinueGeneration',
serverKey: 'enableContinueGeneration',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'fullHeightCodeBlocks',
serverKey: 'fullHeightCodeBlocks',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'systemMessage',
serverKey: 'systemMessage',
type: SyncableParameterType.STRING,
canSync: true
},
{
key: 'showSystemMessage',
serverKey: 'showSystemMessage',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{ key: 'theme', serverKey: 'theme', type: SyncableParameterType.STRING, canSync: true },
{
key: 'copyTextAttachmentsAsPlainText',
serverKey: 'copyTextAttachmentsAsPlainText',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'showRawOutputSwitch',
serverKey: 'showRawOutputSwitch',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'alwaysShowSidebarOnDesktop',
serverKey: 'alwaysShowSidebarOnDesktop',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'showRawModelNames',
serverKey: 'showRawModelNames',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{ key: 'mcpServers', serverKey: 'mcpServers', type: SyncableParameterType.STRING, canSync: true },
{
key: 'agenticMaxTurns',
serverKey: 'agenticMaxTurns',
type: SyncableParameterType.NUMBER,
canSync: true
},
{
key: 'agenticMaxToolPreviewLines',
serverKey: 'agenticMaxToolPreviewLines',
type: SyncableParameterType.NUMBER,
canSync: true
},
{
key: 'showToolCallInProgress',
serverKey: 'showToolCallInProgress',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'alwaysShowAgenticTurns',
serverKey: 'alwaysShowAgenticTurns',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'excludeReasoningFromContext',
serverKey: 'excludeReasoningFromContext',
type: SyncableParameterType.BOOLEAN,
canSync: true
},
{
key: 'sendOnEnter',
serverKey: 'sendOnEnter',
type: SyncableParameterType.BOOLEAN,
canSync: true
}
];
export class ParameterSyncService {
/**
*
@ -298,7 +53,7 @@ export class ParameterSyncService {
// Handle samplers array conversion to string
if (serverParams.samplers && Array.isArray(serverParams.samplers)) {
extracted.samplers = serverParams.samplers.join(';');
extracted[SETTINGS_KEYS.SAMPLERS] = serverParams.samplers.join(';');
}
}

View file

@ -0,0 +1,11 @@
import { ROUTES } from '$lib/constants/routes';
export class RouterService {
static chat(id: string): string {
return `${ROUTES.CHAT}/${id}`;
}
static settings(section: string): string {
return `${ROUTES.SETTINGS}/${section}`;
}
}

View file

@ -413,8 +413,6 @@ class AgenticStore {
const tools = toolsStore.getEnabledToolsForLLM();
if (tools.length === 0) {
console.log('[AgenticStore] No tools available, falling back to standard chat');
return { handled: false };
}

View file

@ -36,7 +36,8 @@ import {
import {
MAX_INACTIVE_CONVERSATION_STATES,
INACTIVE_CONVERSATION_STATE_MAX_AGE_MS,
SYSTEM_MESSAGE_PLACEHOLDER
SYSTEM_MESSAGE_PLACEHOLDER,
TITLE_GENERATION
} from '$lib/constants';
import type {
ChatMessageTimings,
@ -44,7 +45,12 @@ import type {
ChatStreamCallbacks,
ErrorDialogState
} from '$lib/types/chat';
import type { ApiProcessingState, DatabaseMessage, DatabaseMessageExtra } from '$lib/types';
import type {
ApiChatMessageData,
ApiProcessingState,
DatabaseMessage,
DatabaseMessageExtra
} from '$lib/types';
import { ErrorDialogType, MessageRole, MessageType } from '$lib/enums';
interface ConversationStateEntry {
@ -259,7 +265,7 @@ class ChatStore {
}
private isChatLoadingInternal(convId: string): boolean {
return this.chatStreamingStates.has(convId);
return this.chatLoadingStates.has(convId) || this.chatStreamingStates.has(convId);
}
hasPendingMessage(convId: string): boolean {
@ -572,7 +578,11 @@ class ChatStore {
conversationsStore.addMessageToActive(assistantMessage);
await this.streamChatCompletion(
conversationsStore.activeMessages.slice(0, -1),
assistantMessage
assistantMessage,
undefined,
undefined,
undefined,
config().titleGenerationUseLLM && isNewConversation ? content : undefined
);
} catch (error) {
if (isAbortError(error)) {
@ -601,7 +611,8 @@ class ChatStore {
assistantMessage: DatabaseMessage,
onComplete?: (content: string) => Promise<void>,
onError?: (error: Error) => void,
modelOverride?: string | null
modelOverride?: string | null,
firstUserMessageContent?: string
): Promise<void> {
let effectiveModel = modelOverride;
@ -845,6 +856,10 @@ class ChatStore {
perChatOverrides
});
if (agenticResult.handled) {
// Generate LLM based title for new conversations after agentic flow completes
if (firstUserMessageContent) {
await this.generateTitleWithLLM(firstUserMessageContent, streamedContent, convId);
}
// Check if there's a pending steering message to re-send
const pending = agenticStore.consumePendingSteeringMessage(convId);
if (pending) {
@ -894,6 +909,12 @@ class ChatStore {
if (onComplete) await onComplete(content);
if (isRouterMode()) modelsStore.fetchRouterModels().catch(console.error);
// Generate LLM based title for new conversations (avoids stale reference
// issue when user switches conversations while streaming)
if (firstUserMessageContent) {
await this.generateTitleWithLLM(firstUserMessageContent, streamedContent, convId);
}
// Check if there's a pending message queued during streaming
const pending = this.consumePendingMessage(convId);
if (pending) {
@ -921,6 +942,49 @@ class ChatStore {
this.setProcessingState(convId, null);
this.clearPendingMessage(convId);
}
private async generateTitleWithLLM(
userContent: string,
assistantContent: string,
convId: string
): Promise<void> {
const effectiveModel = isRouterMode() && selectedModelName() ? selectedModelName() : undefined;
const configValue = config();
const titlePromptTemplate =
typeof configValue.titleGenerationPrompt === 'string' &&
configValue.titleGenerationPrompt.trim()
? configValue.titleGenerationPrompt
: TITLE_GENERATION.DEFAULT_PROMPT;
const titlePrompt = titlePromptTemplate
.replace('{{USER}}', String(userContent || ''))
.replace('{{ASSISTANT}}', String(assistantContent || ''));
const titleMessage: ApiChatMessageData = {
role: MessageRole.USER,
content: titlePrompt
};
const titleResponse = await ChatService.generateTitle(titleMessage, effectiveModel);
if (!titleResponse) {
return;
}
let cleanTitle = titleResponse.trim();
cleanTitle = cleanTitle
.replace(TITLE_GENERATION.PREFIX_PATTERN, '')
.replace(TITLE_GENERATION.QUOTE_PATTERN, '')
.trim();
if (!cleanTitle || cleanTitle.length < TITLE_GENERATION.MIN_LENGTH) {
const firstLine = userContent.split('\n').find((l) => l.trim().length > 0);
cleanTitle = firstLine ? firstLine.trim() : TITLE_GENERATION.FALLBACK;
}
if (cleanTitle && cleanTitle.length >= TITLE_GENERATION.MIN_LENGTH) {
await conversationsStore.updateConversationName(convId, cleanTitle);
}
}
private async savePartialResponseIfNeeded(convId?: string): Promise<void> {
const conversationId = convId || conversationsStore.activeConversation?.id;
if (!conversationId) return;

View file

@ -44,6 +44,8 @@ import {
MULTIPLE_UNDERSCORE_REGEX,
MCP_DEFAULT_ENABLED_LOCALSTORAGE_KEY
} from '$lib/constants';
import { ROUTES } from '$lib/constants/routes';
import { RouterService } from '$lib/services/router.service';
import { SvelteMap, SvelteSet } from 'svelte/reactivity';
export interface ConversationTreeItem {
@ -260,7 +262,7 @@ class ConversationsStore {
this.activeConversation = conversation;
this.activeMessages = [];
await goto(`#/chat/${conversation.id}`);
await goto(RouterService.chat(conversation.id));
return conversation.id;
}
@ -336,7 +338,7 @@ class ConversationsStore {
if (this.activeConversation && idsToRemove.has(this.activeConversation.id)) {
this.clearActiveConversation();
await goto(`?new_chat=true#/`);
await goto(ROUTES.NEW_CHAT);
}
} else {
// Reparent direct children to deleted conv's parent (or promote to top-level)
@ -352,7 +354,7 @@ class ConversationsStore {
if (this.activeConversation?.id === convId) {
this.clearActiveConversation();
await goto(`?new_chat=true#/`);
await goto(ROUTES.NEW_CHAT);
}
}
} catch (error) {
@ -376,7 +378,7 @@ class ConversationsStore {
toast.success('All conversations deleted');
await goto(`?new_chat=true#/`);
await goto(ROUTES.NEW_CHAT);
} catch (error) {
console.error('Failed to delete all conversations:', error);
toast.error('Failed to delete conversations');
@ -729,7 +731,7 @@ class ConversationsStore {
this.conversations = [newConv, ...this.conversations];
await goto(`#/chat/${newConv.id}`);
await goto(RouterService.chat(newConv.id));
toast.success('Conversation forked');

View file

@ -21,6 +21,7 @@
import { browser } from '$app/environment';
import { base } from '$app/paths';
import { SETTINGS_KEYS } from '$lib/constants';
import { MCPService } from '$lib/services/mcp.service';
import { config, settingsStore } from '$lib/stores/settings.svelte';
import { mcpResourceStore } from '$lib/stores/mcp-resources.svelte';
@ -556,13 +557,13 @@ class MCPStore {
requestTimeoutSeconds: DEFAULT_MCP_CONFIG.requestTimeoutSeconds,
useProxy: serverData.useProxy
};
settingsStore.updateConfig('mcpServers', JSON.stringify([...servers, newServer]));
settingsStore.updateConfig(SETTINGS_KEYS.MCP_SERVERS, JSON.stringify([...servers, newServer]));
}
updateServer(id: string, updates: Partial<MCPServerSettingsEntry>): void {
const servers = this.getServers();
settingsStore.updateConfig(
'mcpServers',
SETTINGS_KEYS.MCP_SERVERS,
JSON.stringify(
servers.map((server) => (server.id === id ? { ...server, ...updates } : server))
)
@ -571,7 +572,10 @@ class MCPStore {
removeServer(id: string): void {
const servers = this.getServers();
settingsStore.updateConfig('mcpServers', JSON.stringify(servers.filter((s) => s.id !== id)));
settingsStore.updateConfig(
SETTINGS_KEYS.MCP_SERVERS,
JSON.stringify(servers.filter((s) => s.id !== id))
);
this.clearHealthCheck(id);
}

View file

@ -10,6 +10,7 @@ import {
MODEL_PROPS_CACHE_MAX_ENTRIES,
FAVORITE_MODELS_LOCALSTORAGE_KEY
} from '$lib/constants';
import { conversationsStore } from '$lib/stores/conversations.svelte';
/**
* modelsStore - Reactive store for model management in both MODEL and ROUTER modes
@ -424,6 +425,103 @@ class ModelsStore {
}
}
/**
* Gets the model name from the last assistant message in the active conversation.
* Iterates backward through messages to find the most recent message with a model.
* Used by both the chat page and settings page to maintain model consistency.
* @returns The model name or null if not found
*/
getModelFromLastAssistantResponse(): string | null {
const messages = conversationsStore.activeMessages;
if (!messages || messages.length === 0) return null;
// Iterate backward to find the last message with a model
for (let i = messages.length - 1; i >= 0; i--) {
if (messages[i].model) {
return messages[i].model;
}
}
return null;
}
/**
* Auto-selects the model from the last assistant response if available and loaded.
* Returns true if a model was selected, false otherwise.
* This is used by the chat page to maintain model consistency across page navigation.
*/
async selectModelFromLastAssistantResponse(): Promise<boolean> {
const lastModel = this.getModelFromLastAssistantResponse();
if (!lastModel) return false;
// Skip if already selected
if (this.selectedModelName === lastModel) return false;
const matchingModel = this.models.find((option) => option.model === lastModel);
if (!matchingModel) return false;
if (!this.isModelLoaded(lastModel)) {
console.log('[modelsStore] last assistant model not loaded:', lastModel);
return false;
}
try {
await this.selectModelById(matchingModel.id);
console.log(`[modelsStore] Automatically selected model: ${lastModel} from last message`);
return true;
} catch (error) {
console.warn('[modelsStore] Failed to automatically select model from last message:', error);
return false;
}
}
/**
* Auto-selects the first available model if none is selected, and fetches its props.
* Prioritizes:
* 1. Model from active conversation's last assistant response (if loaded)
* 2. Model from active conversation's last assistant response (if not loaded)
* 3. First loaded model (not from active conversation)
* 4. First available model
* This is used to ensure default values are populated in settings pages.
*/
async ensureFirstModelSelected(): Promise<void> {
if (this.selectedModelName) return;
// Filter models that are visible in webui
const availableModels = this.models.filter((option) => {
const modelProps = this.getModelProps(option.model);
return modelProps?.webui !== false;
});
if (availableModels.length === 0) return;
// Try to select model from last assistant response first
const lastModel = this.getModelFromLastAssistantResponse();
if (lastModel) {
const lastModelOption = availableModels.find((m) => m.model === lastModel);
if (lastModelOption) {
await this.selectModelById(lastModelOption.id);
if (this.isModelLoaded(lastModel)) {
await this.fetchModelProps(lastModel);
}
return;
}
}
// Try to find a loaded model first
const loadedModel = availableModels.find((m) => this.isModelLoaded(m.model));
if (loadedModel) {
await this.selectModelById(loadedModel.id);
await this.fetchModelProps(loadedModel.model);
return;
}
// Fall back to the first available model
const firstModel = availableModels[0];
await this.selectModelById(firstModel.id);
// Don't fetch props for unloaded models (will fail in ROUTER mode)
}
/**
* Update modalities for a specific model
* Called when a model is loaded or when we need fresh modality data

View file

@ -1,4 +1,6 @@
let _url = $state('#/');
import { SETTINGS_FALLBACK_EXIT_ROUTE } from '$lib/constants';
let _url = $state<string>(SETTINGS_FALLBACK_EXIT_ROUTE);
export const settingsReferrer = {
get url() {

View file

@ -32,9 +32,13 @@
*/
import { browser } from '$app/environment';
import { ColorMode } from '$lib/enums';
import type { SettingsExportType } from '$lib/types';
import { setMode } from 'mode-watcher';
import {
CONFIG_LOCALSTORAGE_KEY,
SETTING_CONFIG_DEFAULT,
SETTINGS_KEYS,
USER_OVERRIDES_LOCALSTORAGE_KEY
} from '$lib/constants';
import { IsMobile } from '$lib/hooks/is-mobile.svelte';
@ -57,7 +61,6 @@ class SettingsStore {
*/
config = $state<SettingsConfigType>({ ...SETTING_CONFIG_DEFAULT });
theme = $state<string>('auto');
isInitialized = $state(false);
userOverrides = $state<Set<string>>(new Set());
@ -99,7 +102,9 @@ class SettingsStore {
initialize() {
try {
this.loadConfig();
this.loadTheme();
this.migrateLegacyTheme();
// Apply the persisted theme from config on initial load
setMode(this.config[SETTINGS_KEYS.THEME] as ColorMode);
this.isInitialized = true;
} catch (error) {
console.error('Failed to initialize settings store:', error);
@ -124,9 +129,9 @@ class SettingsStore {
};
// Default sendOnEnter to false on mobile when the user has no saved preference
if (!('sendOnEnter' in savedVal)) {
if (!(SETTINGS_KEYS.SEND_ON_ENTER in savedVal)) {
if (new IsMobile().current) {
this.config.sendOnEnter = false;
this.config[SETTINGS_KEYS.SEND_ON_ENTER] = false;
}
}
@ -143,12 +148,21 @@ class SettingsStore {
}
/**
* Load theme from localStorage
* Migrate the legacy un-namespaced "theme" localStorage key into config.
* Previously theme was stored separately in localStorage("theme") now it lives
* inside the config object alongside all other settings.
* After migration the legacy key is removed.
*/
private loadTheme() {
private migrateLegacyTheme() {
if (!browser) return;
this.theme = localStorage.getItem('theme') || 'auto';
const legacyTheme = localStorage.getItem('theme');
if (legacyTheme) {
this.config[SETTINGS_KEYS.THEME] = legacyTheme;
localStorage.removeItem('theme');
this.saveConfig();
setMode(legacyTheme as ColorMode);
}
}
/**
*
@ -233,29 +247,13 @@ class SettingsStore {
}
/**
* Update the theme setting
* Update the theme setting.
* @param newTheme - The new theme value
*/
updateTheme(newTheme: string) {
this.theme = newTheme;
this.saveTheme();
}
this.updateConfig(SETTINGS_KEYS.THEME, newTheme);
/**
* Save the current theme to localStorage
*/
private saveTheme() {
if (!browser) return;
try {
if (this.theme === 'auto') {
localStorage.removeItem('theme');
} else {
localStorage.setItem('theme', this.theme);
}
} catch (error) {
console.error('Failed to save theme to localStorage:', error);
}
setMode(newTheme as ColorMode);
}
/**
@ -271,22 +269,26 @@ class SettingsStore {
*/
resetConfig() {
this.config = { ...SETTING_CONFIG_DEFAULT };
this.saveConfig();
}
/**
* Reset theme to auto
* Reset theme to default value.
* Theme is now stored inside the config object.
*/
resetTheme() {
this.theme = 'auto';
this.saveTheme();
this.updateConfig(SETTINGS_KEYS.THEME, SETTING_CONFIG_DEFAULT[SETTINGS_KEYS.THEME]);
setMode(SETTING_CONFIG_DEFAULT[SETTINGS_KEYS.THEME] as ColorMode);
}
/**
* Reset all settings to defaults
* Reset all settings to defaults.
*/
resetAll() {
this.resetConfig();
this.resetTheme();
}
@ -456,10 +458,86 @@ class SettingsStore {
this.saveConfig();
console.log('Cleared all user overrides');
}
/**
*
*
* Import / Export
*
*
*/
/**
* Export all settings as a versioned JSON-compatible object.
* The export captures the full config (excluding sensitive values like API key)
* and user overrides. Sensitive fields are filtered out for security by default.
* @param includeSensitiveData - If true, include sensitive fields (apiKey, MCP server headers) in export
*/
exportSettings(includeSensitiveData: boolean = false): SettingsExportType {
// Build config excluding sensitive data unless user opts in
const configToExport: Record<string, string | number | boolean | undefined> =
includeSensitiveData
? { ...this.config }
: Object.fromEntries(Object.entries(this.config).filter(([key]) => key !== 'apiKey'));
// Handle MCP servers: exclude custom headers unless user opts in
if ('mcpServers' in configToExport && !includeSensitiveData) {
try {
const mcpServers = JSON.parse(configToExport.mcpServers as string) as Array<
Record<string, unknown>
>;
const safeServers = mcpServers.map((server) => {
delete server.headers;
return server;
});
configToExport.mcpServers = JSON.stringify(safeServers);
} catch {
// If parsing fails, just exclude the entire mcpServers field
delete (configToExport as Record<string, unknown>).mcpServers;
}
}
return {
version: 1,
timestamp: Date.now(),
config: configToExport,
userOverrides: Array.from(this.userOverrides)
};
}
/**
* Import settings from a previously exported object.
* Restores config (including theme) and user overrides.
* @param data - The exported settings object
*/
importSettings(data: SettingsExportType): void {
if (!browser) return;
if (!data || !data.config) {
throw new Error('Invalid settings data: missing config');
}
// Restore config (theme is included in config)
this.config = {
...SETTING_CONFIG_DEFAULT,
...data.config
};
// Restore user overrides (derived state — may be stale if server defaults differ)
this.userOverrides = new Set(data.userOverrides ?? []);
// Persist to localStorage
this.saveConfig();
// Apply theme for immediate visual feedback
setMode(this.config[SETTINGS_KEYS.THEME] as ColorMode);
console.log('Settings imported successfully');
}
}
export const settingsStore = new SettingsStore();
export const config = () => settingsStore.config;
export const theme = () => settingsStore.theme;
export const theme = () => settingsStore.config[SETTINGS_KEYS.THEME];
export const isInitialized = () => settingsStore.isInitialized;

View file

@ -76,10 +76,15 @@ export type {
SettingsFieldConfig,
SettingsChatServiceOptions,
SettingsConfigType,
SettingsExportType,
ParameterValue,
ParameterRecord,
ParameterInfo,
SyncableParameter
SyncableParameter,
SettingsEntry,
SettingsSectionTitle,
SettingsSectionEntry,
SettingsSection
} from './settings';
// Common types

View file

@ -1,12 +1,42 @@
import type { SETTING_CONFIG_DEFAULT } from '$lib/constants';
import type { SETTING_CONFIG_DEFAULT, SETTINGS_SECTION_TITLES } from '$lib/constants';
import type { ChatMessagePromptProgress, ChatMessageTimings } from './chat';
import type { OpenAIToolDefinition } from './mcp';
import type { DatabaseMessageExtra } from './database';
import type { ParameterSource, SyncableParameterType, SettingsFieldType } from '$lib/enums';
import type { Icon } from '@lucide/svelte';
import type { Component } from 'svelte';
export type SettingsConfigValue = string | number | boolean | undefined;
/** Section title type derived from registry section titles. */
export type SettingsSectionTitle =
(typeof SETTINGS_SECTION_TITLES)[keyof typeof SETTINGS_SECTION_TITLES];
/** Per-setting metadata — one entry per setting. */
export interface SettingsEntry {
key: string;
label: string;
help: string;
defaultValue: SettingsConfigValue;
type: SettingsFieldType;
section?: string;
options?: Array<{ value: string; label: string; icon: Component }>;
isExperimental?: boolean;
isPositiveInteger?: boolean;
sync?: {
serverKey: string;
paramType: SyncableParameterType;
};
}
/** A settings section with its icon, slug, title, and ordered settings. */
export interface SettingsSectionEntry {
title: SettingsSectionTitle;
slug: string;
icon: Component;
settings: SettingsEntry[];
}
export interface SettingsFieldConfig {
key: string;
label: string;
@ -16,6 +46,14 @@ export interface SettingsFieldConfig {
options?: Array<{ value: string; label: string; icon?: typeof Icon }>;
}
/** Re-exported for backward compatibility. */
export interface SettingsSection {
fields?: SettingsFieldConfig[];
icon: Component;
slug: string;
title: SettingsSectionTitle;
}
export interface SettingsChatServiceOptions {
stream?: boolean;
// Model (required in ROUTER mode, optional in MODEL mode)
@ -94,3 +132,18 @@ export interface SyncableParameter {
type: SyncableParameterType;
canSync: boolean;
}
/**
* Shape of the settings JSON export file.
* Versioned to allow future schema evolution.
*/
export interface SettingsExportType {
/** Export format version — bumped on breaking changes */
version: number;
/** Unix timestamp of export */
timestamp: number;
/** Full settings config (includes theme as a config key) */
config: SettingsConfigType;
/** Keys that differ from server defaults (derived, but persisted for fidelity) */
userOverrides: string[];
}

View file

@ -2,6 +2,7 @@ import { convertPDFToImage, convertPDFToText } from './pdf-processing';
import { isSvgMimeType, svgBase64UrlToPngDataURL } from './svg-to-png';
import { isWebpMimeType, webpBase64UrlToPngDataURL } from './webp-to-png';
import { FileTypeCategory, AttachmentType, SpecialFileType } from '$lib/enums';
import { SETTINGS_KEYS } from '$lib/constants';
import { config, settingsStore } from '$lib/stores/settings.svelte';
import { modelsStore } from '$lib/stores/models.svelte';
import { getFileTypeCategory } from '$lib/utils';
@ -106,7 +107,7 @@ export async function parseFilesToMessageExtras(
console.log('Non-vision model detected: forcing PDF-to-text mode and updating settings');
// Update the setting in localStorage
settingsStore.updateConfig('pdfAsImage', false);
settingsStore.updateConfig(SETTINGS_KEYS.PDF_AS_IMAGE, false);
// Show toast notification to user
toast.warning(

View file

@ -1,6 +1,7 @@
import { isSvgMimeType, svgBase64UrlToPngDataURL } from './svg-to-png';
import { isWebpMimeType, webpBase64UrlToPngDataURL } from './webp-to-png';
import { FileTypeCategory } from '$lib/enums';
import { SETTINGS_KEYS } from '$lib/constants';
import { modelsStore } from '$lib/stores/models.svelte';
import { settingsStore } from '$lib/stores/settings.svelte';
import { toast } from 'svelte-sonner';
@ -104,7 +105,7 @@ export async function processFilesToChatUploaded(
action: {
label: 'Enable PDF as Images',
onClick: () => {
settingsStore.updateConfig('pdfAsImage', true);
settingsStore.updateConfig(SETTINGS_KEYS.PDF_AS_IMAGE, true);
toast.success('PDF parsing as images enabled!', {
duration: 3000
});

View file

@ -7,11 +7,11 @@
import { onMount } from 'svelte';
import { page } from '$app/state';
import { replaceState } from '$app/navigation';
import { APP_NAME } from '$lib/constants';
import { APP_NAME, NEW_CHAT_PARAM } from '$lib/constants';
let qParam = $derived(page.url.searchParams.get('q'));
let modelParam = $derived(page.url.searchParams.get('model'));
let newChatParam = $derived(page.url.searchParams.get('new_chat'));
let newChatParam = $derived(page.url.searchParams.get(NEW_CHAT_PARAM));
// Dialog state for model not available error
let showModelNotAvailable = $state(false);
@ -26,7 +26,7 @@
url.searchParams.delete('q');
url.searchParams.delete('model');
url.searchParams.delete('new_chat');
url.searchParams.delete(NEW_CHAT_PARAM);
replaceState(url.toString(), {});
}

View file

@ -3,13 +3,10 @@
import { page } from '$app/state';
import { afterNavigate } from '$app/navigation';
import { DialogModelNotAvailable } from '$lib/components/app';
import { ROUTES } from '$lib/constants/routes';
import { chatStore, isLoading } from '$lib/stores/chat.svelte';
import {
conversationsStore,
activeConversation,
activeMessages
} from '$lib/stores/conversations.svelte';
import { modelsStore, modelOptions, selectedModelId } from '$lib/stores/models.svelte';
import { conversationsStore, activeConversation } from '$lib/stores/conversations.svelte';
import { modelsStore, modelOptions } from '$lib/stores/models.svelte';
let chatId = $derived(page.params.id);
let currentChatId: string | undefined = undefined;
@ -73,47 +70,9 @@
urlParamsProcessed = true;
}
async function selectModelFromLastAssistantResponse() {
const messages = activeMessages();
if (messages.length === 0) return;
let lastMessageWithModel: DatabaseMessage | undefined;
for (let i = messages.length - 1; i >= 0; i--) {
if (messages[i].model) {
lastMessageWithModel = messages[i];
break;
}
}
if (!lastMessageWithModel) return;
const currentModelId = selectedModelId();
const currentModelName = modelOptions().find((m) => m.id === currentModelId)?.model;
if (currentModelName === lastMessageWithModel.model) {
return;
}
const matchingModel = modelOptions().find(
(option) => option.model === lastMessageWithModel.model
);
if (matchingModel && modelsStore.isModelLoaded(matchingModel.model)) {
try {
await modelsStore.selectModelById(matchingModel.id);
console.log(
`Automatically selected model: ${lastMessageWithModel.model} from last message`
);
} catch (error) {
console.warn('Failed to automatically select model from last message:', error);
}
}
}
afterNavigate(() => {
setTimeout(() => {
selectModelFromLastAssistantResponse();
void modelsStore.selectModelFromLastAssistantResponse();
}, 100);
});
@ -141,7 +100,7 @@
await handleUrlParams();
}
} else {
await goto('#/');
await goto(ROUTES.START);
}
})();
}

View file

@ -2,6 +2,7 @@
import { page } from '$app/stores';
import { goto } from '$app/navigation';
import { ServerErrorSplash } from '$lib/components/app';
import { ROUTES } from '$lib/constants/routes';
let error = $derived($page.error);
let status = $derived($page.status);
@ -17,7 +18,7 @@
function handleRetry() {
// Navigate back to home page after successful API key validation
goto('#/');
goto(ROUTES.START);
}
</script>
@ -60,7 +61,7 @@
</p>
</div>
<button
onclick={() => goto('#/')}
onclick={() => goto(ROUTES.START)}
class="rounded-md bg-primary px-4 py-2 text-primary-foreground hover:bg-primary/90"
>
Go Home

View file

@ -18,6 +18,8 @@
import { isRouterMode, serverStore } from '$lib/stores/server.svelte';
import { config, settingsStore } from '$lib/stores/settings.svelte';
import { ModeWatcher } from 'mode-watcher';
import { ROUTES } from '$lib/constants/routes';
import { RouterService } from '$lib/services/router.service';
import { Toaster } from 'svelte-sonner';
import { modelsStore } from '$lib/stores/models.svelte';
import { mcpStore } from '$lib/stores/mcp.svelte';
@ -53,7 +55,7 @@
const currentId = page.params.id;
if (!currentId) {
goto(`#/chat/${allConvs[direction === 1 ? 0 : allConvs.length - 1].id}`);
goto(RouterService.chat(allConvs[direction === 1 ? 0 : allConvs.length - 1].id));
return;
}
@ -64,9 +66,9 @@
const targetIdx = idx + direction;
if (targetIdx >= 0 && targetIdx < allConvs.length) {
goto(`#/chat/${allConvs[targetIdx].id}`);
goto(RouterService.chat(allConvs[targetIdx].id));
} else {
goto('?new_chat=true#/');
goto(ROUTES.NEW_CHAT);
}
}

View file

@ -4,6 +4,7 @@
import { browser } from '$app/environment';
import { page } from '$app/state';
import { ActionIcon } from '$lib/components/app';
import { SETTINGS_FALLBACK_EXIT_ROUTE } from '$lib/constants';
let { children } = $props();
@ -21,7 +22,7 @@
if (browser && window.history.length > 1 && !prevIsSettings) {
history.back();
} else {
goto('#/');
goto(SETTINGS_FALLBACK_EXIT_ROUTE);
}
}
</script>

View file

@ -0,0 +1,16 @@
<script lang="ts">
import { SettingsChat } from '$lib/components/app/settings';
import { page } from '$app/state';
import { replaceState } from '$app/navigation';
import { RouterService } from '$lib/services';
import { SETTINGS_SECTION_SLUGS } from '$lib/constants';
import { onMount } from 'svelte';
onMount(() => {
if (!page.params.section) {
replaceState(RouterService.settings(SETTINGS_SECTION_SLUGS.GENERAL), {});
}
});
</script>
<SettingsChat initialSection={(page.params as Record<string, string | undefined>).section} />

View file

@ -1,6 +0,0 @@
<script lang="ts">
import { SettingsChat } from '$lib/components/app/settings';
import { page } from '$app/state';
</script>
<SettingsChat initialSection={(page.params as Record<string, string | undefined>).section} />

View file

@ -96,6 +96,7 @@ export default defineConfig({
'/props': 'http://localhost:8080',
'/models': 'http://localhost:8080',
'/tools': 'http://localhost:8080',
'/slots': 'http://localhost:8080',
'/cors-proxy': 'http://localhost:8080'
},
headers: {