Compare commits

..

16 Commits

Author SHA1 Message Date
Yiwei Shao 1a03cf47f6 hexagon: hmx flash attention (#22347)
* hmx: extract shared interleave headers and unify matmul batched

* hmx: add HMX-accelerated flash attention for prefill

* hmx: replace asm wrappers with Q6_ intrinsics in hmx-utils.h

Switches three single-instruction helpers from inline asm to the matching
Q6_ intrinsics, matching the style established by aizip f8737609a and used
by the upstream PR #21554 hmx-matmul-ops.c rewrite:

  hmx_set_output_scales       asm "bias=mxmem2"  -> Q6_bias_mxmem2_A
  hmx_load_tile_pair_fp16     asm packet         -> Q6_activation_hf_mxmem_RR
                                                    + Q6_weight_hf_mxmem_RR
  hmx_consume_accumulator_fp16 asm "mxmem=acc"   -> Q6_mxmem_AR_after_hf

hmx_load_tiles_fp16 stays on inline asm: it uses ":deep" activation
streaming, and the mixed Q6_activation_hf_mxmem_RR_deep + non-deep
Q6_weight_hf_mxmem_RR pair fails the HMX backend constraint check
("activate weight pair (1) exceeds limit (1)"). The asm bundle keeps
both halves in one VLIW packet and avoids the diagnostic.

Functionally equivalent — same instructions emitted; the Q6_ intrinsics
just give the compiler more visibility for scheduling.

* hmx: drop the duplicate interleave_fp16_weight_chunk_to_tiles

* hmx:  apply upstream optimization to hmx-flash-attn-ops.c
apply restrict, __builtin_assume, and pointer accumulation to the three HMX workers (qk_dot, o_update, o_norm) and the matching inline HMX loops in op_hmx_flash_attn_ext.

* hmx: unify interleave helper

* hmx: multi-thread Q load / O store and enable prefill FA dispatch

Extract inline Q-load and O-store loops into worker_pool-parallel helpers
(fa_phase_q_load, fa_phase_o_store) so HVX threads split the F32↔F16
conversion work across row ranges.  Also relax the softmax threading
gate from n_row_vec_cnt >= n_threads to >= 2, which was unnecessarily
forcing single-thread fallback when n_rows_g < 512.

On the dispatch side, remove the ne[2] != 1 guard that blocked multi-head
(prefill) FA from reaching the HTP backend — GQA is already handled
internally by both the HMX and HVX flash-attention paths.

* hmx: relax matmul pipeline gate to cover k > n shapes (e.g. FFN_down)

* hmx: optimize FA softmax mask phase (no-ALiBi fast path + GQA dedup)

* hmx: Add an asm memory clobber at the phase boundary to prevent reorder bug

* [experimental]: fp16 softmax (EXP2_HF) to accelerate fa

Bake log2(e) into qk_scale and use hvx_exp2_hf directly for P and m_diff
(base-2 consistent, matches htp-ops-lib). ~22 ALU ops for 64 lanes vs
~44 for the F32 round-trip path.

* hmx flash-attn: refine cost model coefficients based on profiling data

* hmx flash-attn: replace asm clobber with targeted volatile reads on vtcm_d_tiles

* hmx flash-attn: fix prefill correctness (dst indexing, softmax reduce, V stride)

* hmx flash-attn: fix p_tiles dual-tile OOB race; enable MT + pipeline

* hmx flash-attn: preserve additive mask bias in no-ALiBi fast path

The no-ALiBi fast path (max_bias==0) was skipping mask add entirely on
the assumption that mask values are only {0, -inf}.  This is wrong when
the mask carries additive positional bias — those terms were silently
dropped.  Keep the slope-mul skip (slope≡1.0) but add mask back so the
bias survives; vmux still clamps below -16 to -inf.

Also add HMX FA coverage to test-backend-ops: prefill shapes (nb=64,
nb=32) × {mask on/off} × {ALiBi on/off} × {softcap on/off}, F16 KV,
hs ∈ {64, 128}.

* hmx: fix softcap+EXP2_HF interaction, tighten matmul pipeline gate, add FA tests

- flash-attn: when EXP2_HF is on AND logit_softcap is active, fold
  log2(e) into the post-tanh multiplier (v_cap) instead of pre-baking
  it into qk_scale.  Pre-baking shifted the tanh knee from x≈c to
  x≈c/log2(e) and produced numerically wrong softcapped outputs
  whenever both knobs were enabled.
- flash-attn softmax (fa_softmax_thread): replace the union+memcpy
  scalar extract pattern with HVX vmux-based per-row accumulators on
  rowmax/rowsum.  Add hvx_vec_get_f16 helper in hvx-base.h.  Functional
  parity, less scalar code, clearer hf/qf16 lane-format contract.
- matmul (hmx_mat_mul_permuted_qk_0_d16a32): pick pipeline vs sequential
  layout based on whether the chunker actually yields >=2 n-chunks,
  instead of the static (m>=128 && n>=256) gate.  Avoids paying for
  output double-buffer + worker dispatch when there is no HMX/HVX
  overlap to gain (e.g. shapes that collapse to one n-chunk).
- tests: add HMX flash-attention coverage over the
  {mask, ALiBi (max_bias), logit_softcap} cross-product for the prefill
  path — head_dim 64/128, GQA 4×4, kv=512/nb=64 plus a kv=113/nb=32
  non-aligned case.

* [Help Wanted]: refactor D matrix computation into separate function for clarity and maintainability

* format code

* hexagon: looks like -O3 is causing issues with the large code base, switch to -O2 and -flto instead

* hexagon: use hex_ prefix for swap_ptr

* hexagon: move vtcm_seq_alloc into vtcm-utils.h

More vtcm allocator updates are coming so it makes sense to start the separate hdr for it.

* hmx-utils: add hmx_prefix for layout converters

* hmx-mm: move main hmx_mm functions to the end, remove unused fwd decls, etc

* hmx-mm: remove unused qweight_fetch_task_state_t and minor alignment fixes

* hmx-fa: minor alignment fixes

* hmx-fa: move hmx_flash_atten into hmx-ops.h

* hmx-fa: remove redundant workpool pointer in the hmx_fa_ctx, plus minor alignment updates

* hmx-fa: minor alignment and simplifications

* hexagon: move FA_EXP_F16 option to hostside CMake file

* hmx-fa: use hvx_vec_splat_f16 instead of fp16_to_bits

* hmx-fa: add hvx_splat_u16/u8 and use that in the fa instead custom hvx_fill

* hmx-fa: some more alignment updates in the core fa function

* hmx-fa: keep slopes in vtcm in fp16

Saves malloc/free and removes the need for float -> fp16 downcast on every use.

* hexagon: consistent noinline usage (after static)

* hex-hmx: consistent use FARF_HIGH to enable debug output

* hmx-utils: no need for always_inline attr

* hex-hmx: consistent noinline usage (static noinline ...)

* hex-hmx: simplify init_col_scales

* hexagon: fix editorconfig errors

* hmx-mm: minor alignment fixes

---------

Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
2026-05-01 20:29:13 -07:00
ddh0 b97ebdc98f llama-quant : fix --tensor-type when default qtype is overriden (#22572)
fix #22544 (my fault!)

Credit to @Anai-Guo, ref #22559 - since that one was closed due to the
new contributor policy I am taking the liberty of re-submitting that PR
here.
2026-05-01 19:55:55 +02:00
Aparna M P 2098fd6169 hexagon: enable non-contiguous row tensor support for unary ops (#22574) 2026-05-01 10:09:23 -07:00
Aleksander Grygier ab6120cde5 webui: Spring Cleaning Refactor v1 (#22505)
* wip: server_tools

* feat: Integrate with `/tools` endpoint

* feat: Builtin + MCP + JSON Schema Tools WIP

* refactor

* displayName -> display_name

* snake_case everywhere

* rm redundant field

* feat: Improvements

* chore: update webui build output

* refactor: Updates after server updates

* chore: update webui build output

* change arg to --tools all

* feat: UI improvements

* chore: update webui build output

* add readme mention

* llama-gen-docs

* chore: update webui build output

* chore: update webui build output

* chore: update webui build output

* feat: Reorganize settings sections

* feat: Separate dialogs for MCP Servers Settings and Import/Export

* feat: WIP

* feat: WIP

* feat: WIP

* feat: WIP

* feat: WIP

* feat: WIP

* WIP on allozaur/20677-webui-server-tools

* feat: UI improvements

* chore: Update package lock

* chore: Run `npm audit fix`

* feat: UI WIP

* feat: UI

* refactor: Desktop Icon Strip DRY

* feat: Cleaner rendering and transition for ChatScreen

* feat: UI improvements

* feat: UI improvement

* feat: Remove MCP Server "enable" switch from Tools submenu

* chore: Run `npm audit fix`

* feat: WIP

* feat: Logic improvements

* refactor: Cleanup

* refactor: DRY

* test: Fix Chat Sidebar UI Tests

* chore: Update package lock

* refactor: Cleanup

* feat: Chat Message Action Card with Continue and Permission flow implementations

* feat: Add agentic steering messages, draft messages and improve chat UX

* fix: Search results UI

* test: Fix unit test

* feat: UI/UX improvements

* refactor: Simplify `useToolsPanel` access in components

* feat: Implement Processing Info Context API

* feat: Implement 'Go back to chat' functionality for settings

* feat: Enhance MCP Server management in Chat Form Attachments

* style: Minor UI and branding adjustments

* chore: Update webui static build output

* chore: Formatting, linting & type checks

* feat: Draft messages logic

* feat: UI improvements

* feat: Steering Messages improvements

* refactor: Cleanup

* refactor: Cleanup

* feat: Improve UI

* refactor: Settings navigation hook

* refactor: DRY code

* refactor: DRY ChatMessageUser UI components

* refactor: Desktop Icon Strip DRY

* refactor: Tools & permissions

* fix: Navigation condition

* refactor: Cleanup

* refactor: Cleanup

* refactor: Cleanup

* fix: preserve reasoning_content in agentic flow

* refactor: Storybook cleanup

* refactor: isInViewport util function

* refactor: Rename globally `onClick` to `onclick`

* chore: `npm audit fix`

* refactor: Action Icon usage

* refactor: Naming

* refactor: JS in `class` directive

* refactor: Chat components cleanup WIP

* refactor: Components structure

* refactor: Cleanup WIP

* feat: New ChatAttachmentsPreview component

* feat: UI improvements

* feat: UI improvements

* refactor: Cleanup

* refactor: ChatAttachmentsPreview UI/UX

* refactor: Remove dead code

* refactor: Cleanup

* fix: Model Name aliases displaying

* feat: Shortcut improvements

* refactor: Chat Message

* feat: Move Import/Export to settings

* refactor: Cleanup

* refactor: Cleanup

* refactor: Cleanup

* refactor: Cleanup

---------

Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
2026-05-01 18:36:29 +02:00
Masashi Yoshimura c3c1505392 ggml-webgpu: Fix vectorized handling in mul-mat and mul-mat-id (#22578)
* Fix vectorized condition of mul-mat-fast pipeline and add vectorized variant to mul-mat-id

* Apply suggestion from @CISC

Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>

---------

Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
2026-05-01 07:55:01 -07:00
Jeff Bolz 05e141a6b3 vulkan: Support asymmetric FA in coopmat2 path (#21753)
* vulkan: Support asymmetric FA in coopmat2 path

There has been some recent interest/experimentation with mixed quantization
types for FA. I had originally designed the cm2 FA shader with this in mind
(because I didn't realize it wasn't supported at the time!), this change
adds the missing pieces and enables it.

Also support Q1_0 since people have been trying that out (seems crazy, but
who knows).

We should be able to do similar things in the coopmat1/scalar path, but
there's another change open against the scalar path and I don't want to
conflict.

* reorder cases
2026-05-01 15:28:32 +02:00
Chen Yuan aab68217b7 ggml-webgpu: add the upscale shader (#22419)
* shader(upscale): add the upscale shader with nearest, bilinear and bicubic implementations

* shader(upscale): use macro
2026-04-30 22:22:18 -07:00
Masashi Yoshimura a95a11e5b8 ggml-webgpu: Improve performance of mat-vec and mat-mat for MUL_MAT_ID (#22464)
* Add mat-vec fast path of MUL_MAT_ID.

* Add shared accumulation vec logic and the other types supports.

* Add i-quant mat-mat for MUL_MAT_ID and fix some parts

* Remove n_experts from shader_lib_context.
2026-04-30 14:19:10 -07:00
Reese Levine 5cbfb18075 Update llama-mmap to use ftello/fseeko (#22497)
* Update llama-mmap to work with 32-bit wasm and >2GB models

* Update to gguf.cpp style
2026-04-30 14:17:52 -07:00
Adrien Gallouët beb42fffa4 common : check for null getpwuid in hf-cache (#22550)
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
2026-04-30 21:32:41 +02:00
Ruben Ortlam 660b1b4bdc vulkan: add get/set tensor 2d functions (#22514)
* vulkan: add get/set_tensor_2d functions

* fix backend interface comments

* Update ggml/src/ggml-metal/ggml-metal.cpp

Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
2026-04-30 17:37:13 +02:00
Ben Guidarelli c20c44514a spec: fix argument typo (#22552) 2026-04-30 17:32:32 +03:00
Sigbjørn Skjæret 6118c043b1 ci : bump ty to 0.0.33 (#22535)
* bump ty to 0.0.33

* update typings
2026-04-30 16:15:54 +03:00
Adrien Gallouët 5f0ab726f7 vendor : update cpp-httplib to 0.43.2 (#22548)
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
2026-04-30 15:04:39 +02:00
Johannes Gäßler e82aaf2587 CUDA: fix tile FA kernel on Pascal (#22541) 2026-04-30 13:04:50 +02:00
Georgi Gerganov 27aef3dd91 scripts : add wc2wt.sh - create worktree from current HEAD (#22513)
* scripts : add wc2wt.sh - create worktree from current HEAD

Add a script to create a git worktree on a new branch from the current
HEAD. Similar to pr2wt.sh but for local development branches instead of
PRs.

Usage:
  ./scripts/wc2wt.sh gg/new-feature
  ./scripts/wc2wt.sh gg/new-feature "bash -l"

Assisted-by: llama.cpp:local pi

* cont : no need to try to delete the branch
2026-04-30 09:20:26 +03:00
217 changed files with 12529 additions and 9562 deletions
+1 -1
View File
@@ -31,7 +31,7 @@ jobs:
uses: actions/setup-python@v6
with:
python-version: "3.11"
pip-install: -r requirements/requirements-all.txt ty==0.0.26
pip-install: -r requirements/requirements-all.txt ty==0.0.33
# - name: Type-check with Pyright
# uses: jakebailey/pyright-action@v2
# with:
+1 -1
View File
@@ -3499,7 +3499,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_LOOKUP, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPEC_DRAFT_N_MIN"));
add_opt(common_arg(
{"--spec--draft-p-split", "--draft-p-split"}, "P",
{"--spec-draft-p-split", "--draft-p-split"}, "P",
string_format("speculative decoding split probability (default: %.2f)", (double)params.speculative.draft.p_split),
[](common_params & params, const std::string & value) {
params.speculative.draft.p_split = std::stof(value);
+1 -1
View File
@@ -57,7 +57,7 @@ static fs::path get_cache_directory() {
#ifndef _WIN32
const struct passwd * pw = getpwuid(getuid());
if (pw->pw_dir && *pw->pw_dir) {
if (pw && pw->pw_dir && *pw->pw_dir) {
return fs::path(pw->pw_dir) / ".cache" / "huggingface" / "hub";
}
#endif
+1 -1
View File
@@ -6658,7 +6658,7 @@ class BertModel(TextModel):
tokens: list[bytes] = [f"[PAD{i}]".encode("utf-8") for i in range(vocab_size)]
scores: list[float] = [-10000.0] * vocab_size
toktypes: list[int] = [SentencePieceTokenTypes.UNUSED] * vocab_size # ty: ignore[invalid-assignment]
toktypes: list[int] = [SentencePieceTokenTypes.UNUSED] * vocab_size
if isinstance(tokenizer, SentencePieceProcessor):
for token_id in range(tokenizer.vocab_size()):
+1 -1
View File
@@ -2100,8 +2100,8 @@ static const ggml_backend_i ggml_backend_meta_i = {
/* .free = */ ggml_backend_meta_free,
/* .set_tensor_async = */ ggml_backend_meta_set_tensor_async,
/* .get_tensor_async = */ ggml_backend_meta_get_tensor_async,
/* .get_tensor_2d_async = */ nullptr,
/* .set_tensor_2d_async = */ nullptr,
/* .get_tensor_2d_async = */ nullptr,
/* .cpy_tensor_async = */ nullptr,
/* .synchronize = */ ggml_backend_meta_synchronize,
/* .graph_plan_create = */ nullptr,
+2 -2
View File
@@ -262,9 +262,9 @@ static struct ggml_backend_i blas_backend_i = {
/* .get_name = */ ggml_backend_blas_get_name,
/* .free = */ ggml_backend_blas_free,
/* .set_tensor_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .set_tensor_2d_async = */ NULL,
/* .get_tensor_async = */ NULL,
/* .set_tensor_2d_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .cpy_tensor_async = */ NULL,
/* .synchronize = */ NULL,
/* .graph_plan_create = */ NULL,
+1 -1
View File
@@ -2746,8 +2746,8 @@ static const ggml_backend_i ggml_backend_cann_interface = {
/* .free = */ ggml_backend_cann_free,
/* .set_tensor_async = */ ggml_backend_cann_set_tensor_async,
/* .get_tensor_async = */ ggml_backend_cann_get_tensor_async,
/* .get_tensor_2d_async = */ NULL,
/* .set_tensor_2d_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .cpy_tensor_async = */ ggml_backend_cann_cpy_tensor_async,
/* .synchronize = */ ggml_backend_cann_synchronize,
/* .graph_plan_create = */ NULL,
+1 -1
View File
@@ -195,8 +195,8 @@ static const struct ggml_backend_i ggml_backend_cpu_i = {
/* .free = */ ggml_backend_cpu_free,
/* .set_tensor_async = */ NULL,
/* .get_tensor_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .set_tensor_2d_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .cpy_tensor_async = */ NULL,
/* .synchronize = */ NULL,
/* .graph_plan_create = */ ggml_backend_cpu_graph_plan_create,
+16 -5
View File
@@ -68,7 +68,7 @@ static constexpr __host__ __device__ uint32_t ggml_cuda_fattn_tile_get_config_nv
GGML_CUDA_FATTN_TILE_CONFIG_CASE(256, 256, 16, 256, 2, 64, 64)
GGML_CUDA_FATTN_TILE_CONFIG_CASE(256, 256, 32, 256, 2, 64, 64)
GGML_CUDA_FATTN_TILE_CONFIG_CASE(320, 256, 32, 256, 2, 64, 64)
GGML_CUDA_FATTN_TILE_CONFIG_CASE(320, 256, 16, 256, 2, 64, 64)
GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 4, 128, 2, 64, 64)
GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 8, 256, 2, 64, 64)
@@ -130,7 +130,7 @@ static constexpr __host__ __device__ uint32_t ggml_cuda_fattn_tile_get_config_nv
GGML_CUDA_FATTN_TILE_CONFIG_CASE(256, 256, 16, 256, 2, 32, 128)
GGML_CUDA_FATTN_TILE_CONFIG_CASE(256, 256, 32, 256, 2, 32, 64)
GGML_CUDA_FATTN_TILE_CONFIG_CASE(320, 256, 32, 256, 2, 32, 64)
GGML_CUDA_FATTN_TILE_CONFIG_CASE(320, 256, 16, 256, 2, 32, 64)
GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 4, 128, 2, 32, 64)
GGML_CUDA_FATTN_TILE_CONFIG_CASE(512, 512, 8, 256, 2, 32, 64)
@@ -1124,7 +1124,7 @@ static void launch_fattn_tile_switch_ncols1(ggml_backend_cuda_context & ctx, ggm
constexpr size_t nbytes_shared = 0;
#ifdef GGML_USE_HIP
if constexpr (DV <= 128) {
if constexpr (DKQ <= 128) {
if (Q->ne[1] > 32/ncols2) {
constexpr int cols_per_block = 64;
const int nwarps = ggml_cuda_fattn_tile_get_nthreads (DKQ, DV, cols_per_block, cc) / warp_size;
@@ -1138,7 +1138,7 @@ static void launch_fattn_tile_switch_ncols1(ggml_backend_cuda_context & ctx, ggm
#endif // GGML_USE_HIP
#ifndef GGML_USE_HIP
if constexpr (DV <= 256)
if constexpr (DKQ <= 256)
#endif // GGML_USE_HIP
{
if (Q->ne[1] > 16/ncols2) {
@@ -1220,11 +1220,22 @@ static void launch_fattn_tile_switch_ncols2(ggml_backend_cuda_context & ctx, ggm
const int gqa_limit = nvidia && gqa_ratio <= 4 && DV <= 256 ? 16 : INT_MAX;
const bool use_gqa_opt = mask && max_bias == 0.0f && Q->ne[1] <= gqa_limit && K->ne[1] % FATTN_KQ_STRIDE == 0;
if constexpr (DKQ == 320) { // Mistral Small 4
if constexpr (DKQ == 320) {
// This branch is only used for Mistral Small 4 which has a GQA ratio of 32.
// On AMD, simply use that GQA ratio with 32 columns / block since we always have enough SRAM.
// On NVIDIA however, the tile kernel is only used for GPUs that can't use the mma kernel (Pascal and older).
// Therefore, use a GQA ratio of 16 with 16 columns / block to stay below 48 kiB of SRAM / block.
#ifdef GGML_USE_HIP
if (use_gqa_opt && gqa_ratio % 32 == 0) {
launch_fattn_tile_switch_ncols1<DKQ, DV, 32, use_logit_softcap>(ctx, dst);
return;
}
#else
if (use_gqa_opt && gqa_ratio % 16 == 0) {
launch_fattn_tile_switch_ncols1<DKQ, DV, 16, use_logit_softcap>(ctx, dst);
return;
}
#endif // GGML_USE_HIP
GGML_ABORT("flash-attn tile (320/256): expected GQA ratio multiple of 32");
}
+2 -2
View File
@@ -4588,8 +4588,8 @@ static const ggml_backend_i ggml_backend_cuda_interface = {
/* .free = */ ggml_backend_cuda_free,
/* .set_tensor_async = */ ggml_backend_cuda_set_tensor_async,
/* .get_tensor_async = */ ggml_backend_cuda_get_tensor_async,
/* .get_tensor_2d_async = */ ggml_backend_cuda_set_tensor_2d_async,
/* .set_tensor_2d_async = */ ggml_backend_cuda_get_tensor_2d_async,
/* .set_tensor_2d_async = */ ggml_backend_cuda_set_tensor_2d_async,
/* .get_tensor_2d_async = */ ggml_backend_cuda_get_tensor_2d_async,
/* .cpy_tensor_async = */ ggml_backend_cuda_cpy_tensor_async,
/* .synchronize = */ ggml_backend_cuda_synchronize,
/* .graph_plan_create = */ NULL,
+2 -1
View File
@@ -22,7 +22,8 @@ message(STATUS "hexagon: using ${HEXAGON_SDK_ROOT} and ${HEXAGON_TOOLS_ROOT} for
include(${HEXAGON_SDK_ROOT}/build/cmake/hexagon_fun.cmake)
include(ExternalProject)
option(GGML_HEXAGON_HTP_DEBUG "ggml-hexagon: enable HTP debug output" OFF)
option(GGML_HEXAGON_HTP_DEBUG "ggml-hexagon: enable HTP debug output" OFF)
option(GGML_HEXAGON_FA_EXP2_HF "ggml-hexagon: use FP16 exp2 polynomial in FA softmax instead of F32 exp round-trip" OFF)
set(GGML_HEXAGON_HTP_CERT "$ENV{HEXAGON_HTP_CERT}" CACHE PATH "ggml-hexagon: enable HTP library signing using certificate")
set(GGML_HEXAGON_FP32_QUANTIZE_GROUP_SIZE 128 CACHE STRING "ggml-hexagon: quantize group size (32, 64, or 128)")
+4 -5
View File
@@ -2254,8 +2254,7 @@ static bool ggml_hexagon_supported_flash_attn_ext(const struct ggml_hexagon_sess
return false;
}
if (dst->ne[2] != 1 || dst->ne[3] != 1) {
// FA during prompt still needs work
if (dst->ne[3] != 1) {
return false;
}
@@ -2421,8 +2420,8 @@ static bool ggml_hexagon_supported_unary(const struct ggml_hexagon_session * ses
return false;
}
// TODO: add support for non-contigiuos tensors
if (!ggml_is_contiguous(src0) || !ggml_is_contiguous(dst)) {
// TODO: add support for non-contiguous elements within a row
if (!ggml_is_contiguous_rows(src0) || !ggml_is_contiguous_rows(dst)) {
return false;
}
@@ -3036,8 +3035,8 @@ static struct ggml_backend_i hexagon_backend_i = {
/* .free = */ ggml_backend_hexagon_free,
/* .set_tensor_async = */ NULL,
/* .get_tensor_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .set_tensor_2d_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .cpy_tensor_async = */ NULL,
/* .synchronize = */ ggml_backend_hexagon_synchronize,
/* .graph_plan_create = */ NULL,
+7
View File
@@ -44,6 +44,11 @@ target_compile_definitions(${HTP_LIB} PRIVATE
$<IF:$<BOOL:${HEXAGON_HTP_DEBUG}>,FARF_HIGH=1,>
FP32_QUANTIZE_GROUP_SIZE=${GGML_HEXAGON_FP32_QUANTIZE_GROUP_SIZE})
if (GGML_HEXAGON_FA_EXP2_HF)
message(STATUS "ggml-htp: HMX_FA_USE_EXP2_HF=1 (use FP16 exp2 polynomial in FA softmax)")
target_compile_definitions(${HTP_LIB} PRIVATE HMX_FA_USE_EXP2_HF=1)
endif()
# HMX acceleration: available on v73+ architectures
set(HTP_HMX_VERSIONS v73 v75 v79 v81)
list(FIND HTP_HMX_VERSIONS ${DSP_VERSION} _hmx_idx)
@@ -52,11 +57,13 @@ if (_hmx_idx GREATER_EQUAL 0)
target_sources(${HTP_LIB} PRIVATE
hmx-queue.c
hmx-matmul-ops.c
hmx-flash-attn-ops.c
)
# -mhmx enables HMX instruction set (needed by files that include hmx-utils.h)
set_source_files_properties(
hmx-matmul-ops.c
hmx-flash-attn-ops.c
PROPERTIES COMPILE_OPTIONS "-mhmx"
)
@@ -138,15 +138,15 @@ set(CMAKE_SHARED_LIBRARY_SONAME_C_FLAG "-Wl,-soname,")
set(CMAKE_SHARED_LIBRARY_SONAME_CXX_FLAG "-Wl,-soname,")
#Compiler Options
set(COMMON_FLAGS "-mcpu=hexagon${V_ARCH} -m${V_ARCH} -mhvx=${V_ARCH} -fvectorize -Wall -Werror -fno-zero-initialized-in-bss -G0 -fdata-sections -fpic ${XQF_ARGS}")
set(COMMON_FLAGS "-mcpu=hexagon${V_ARCH} -m${V_ARCH} -mhvx=${V_ARCH} -fvectorize -flto -Wall -Werror -fno-zero-initialized-in-bss -G0 -fdata-sections -fpic ${XQF_ARGS}")
set(CMAKE_CXX_FLAGS_DEBUG "${COMMON_FLAGS} -O0 -D_DEBUG -g")
set(CMAKE_CXX_FLAGS_RELWITHDEBINFO "${COMMON_FLAGS} -O3 -g")
set(CMAKE_CXX_FLAGS_RELEASE "${COMMON_FLAGS} -O3")
set(CMAKE_CXX_FLAGS_RELWITHDEBINFO "${COMMON_FLAGS} -O2 -g")
set(CMAKE_CXX_FLAGS_RELEASE "${COMMON_FLAGS} -O2")
set(CMAKE_C_FLAGS_DEBUG "${COMMON_FLAGS} -O0 -D_DEBUG -g")
set(CMAKE_C_FLAGS_RELWITHDEBINFO "${COMMON_FLAGS} -O3 -g")
set(CMAKE_C_FLAGS_RELEASE "${COMMON_FLAGS} -O3")
set(CMAKE_C_FLAGS_RELWITHDEBINFO "${COMMON_FLAGS} -O2 -g")
set(CMAKE_C_FLAGS_RELEASE "${COMMON_FLAGS} -O2")
set(CMAKE_ASM_FLAGS_DEBUG "${COMMON_FLAGS} ${CMAKE_CXX_FLAGS_DEBUG}")
set(CMAKE_ASM_FLAGS_RELEASE "${COMMON_FLAGS} ${CMAKE_CXX_FLAGS_RELEASE}")
+13 -1
View File
@@ -17,13 +17,14 @@
#include "htp-ctx.h"
#include "htp-ops.h"
#include "htp-ops.h"
#include "hmx-ops.h"
// Must be multiple of 32
#define FLASH_ATTN_BLOCK_SIZE (32 * 2)
// This is a bit of a hack because the compiler is strugling to properly inline
// the default hvx_vec_f32_to_f16 with output into the local array.
static void __attribute__((noinline)) hvx_vec_f32_to_f16_a(void *ptr, HVX_Vector v0, HVX_Vector v1)
static __attribute__((noinline)) void hvx_vec_f32_to_f16_a(void *ptr, HVX_Vector v0, HVX_Vector v1)
{
*(HVX_Vector *) ptr = hvx_vec_f32_to_f16(v0, v1);
}
@@ -621,6 +622,17 @@ int op_flash_attn_ext(struct htp_ops_context * octx) {
return HTP_STATUS_NO_SUPPORT;
}
#ifdef HTP_HAS_HMX
// HMX path: prefill (neq1 >= 32), head_dim multiple of 32, F16 KV
if (k->type == HTP_TYPE_F16 && v->type == HTP_TYPE_F16 && k->ne[0] % 32 == 0 && q->ne[1] >= 32) {
int ret = hmx_flash_attn_ext(octx);
if (ret == HTP_STATUS_OK) {
return ret;
}
// VTCM too small or other failure -> fall through to HVX path
}
#endif
struct htp_fa_context factx;
factx.octx = octx;
+6
View File
@@ -74,6 +74,12 @@ static inline size_t hex_smax(size_t a, size_t b) {
return a > b ? a : b;
}
static inline void hex_swap_ptr(void ** p1, void ** p2) {
void * t = *p1;
*p1 = *p2;
*p2 = t;
}
static inline void hex_l2fetch(const void * p, uint32_t width, uint32_t stride, uint32_t height) {
const uint64_t control = Q6_P_combine_RR(stride, Q6_R_combine_RlRl(width, height));
Q6_l2fetch_AP((void *) p, control);
File diff suppressed because it is too large Load Diff
File diff suppressed because it is too large Load Diff
+3
View File
@@ -61,6 +61,9 @@ int hmx_mat_mul_permuted_qk_0_d16a32(struct htp_context *ctx,
int m, int k, int n,
int weight_type);
// HMX flash attention
int hmx_flash_attn_ext(struct htp_ops_context * octx);
#ifdef __cplusplus
}
#endif
+181 -11
View File
@@ -4,6 +4,9 @@
#ifndef HMX_UTILS_H
#define HMX_UTILS_H
#include "hvx-base.h"
#include <assert.h>
#include <hexagon_types.h>
#include <stddef.h>
@@ -12,21 +15,188 @@
#define HMX_FP16_TILE_N_ELMS 1024
#define HMX_FP16_TILE_SIZE 2048
#define HMX_INLINE_ALWAYS inline __attribute__((unused, always_inline))
// Initialise aligned 256-byte area with scale vector + zero padding.
static HMX_INLINE_ALWAYS void hmx_init_column_scales(void *out_scales, HVX_Vector v_scale) {
HVX_Vector *pv = (HVX_Vector *)out_scales;
*pv++ = v_scale;
*pv = Q6_V_vzero();
static inline void hmx_init_column_scales(void *out_scales, HVX_Vector v_scale) {
volatile HVX_Vector *pv = (HVX_Vector *) out_scales;
pv[0] = v_scale;
pv[1] = Q6_V_vzero();
}
// --- VTCM sequential allocator (from htp-ops-lib/include/dsp/vtcm_mgr.h) ---
// --- Shared scatter offsets and interleave helper ---
static inline uint8_t *vtcm_seq_alloc(uint8_t **vtcm_ptr, size_t size) {
uint8_t *p = *vtcm_ptr;
*vtcm_ptr += size;
return p;
// vscatter offsets for fused dequant+transpose: write K-values directly to [K][N] tile.
// word[i] = i*128 maps K-row-pair i to byte offset i*128.
// Column offset (n*4) is added at runtime. Entries 0..15 cover one tile (region 2047);
// entries 16..31 cover the next adjacent tile (region 4095) — pick region size at the
// call site to scatter into one tile (masked) or two contiguous tiles (unmasked).
static const int32_t hmx_transpose_scatter_offsets[32] __attribute__((aligned(VLEN))) = {
0 * 128, 1 * 128, 2 * 128, 3 * 128, 4 * 128, 5 * 128, 6 * 128, 7 * 128, 8 * 128, 9 * 128, 10 * 128,
11 * 128, 12 * 128, 13 * 128, 14 * 128, 15 * 128, 16 * 128, 17 * 128, 18 * 128, 19 * 128, 20 * 128, 21 * 128,
22 * 128, 23 * 128, 24 * 128, 25 * 128, 26 * 128, 27 * 128, 28 * 128, 29 * 128, 30 * 128, 31 * 128,
};
// Scatter row-major FP16 data (in VTCM scratch) into transposed [K][N] tiles.
// vtcm_src: [n_cols][src_stride] row-major fp16 (only first k elements per row are used)
// vtcm_dst: [n_col_tiles][n_k_tiles][HMX_FP16_TILE_N_ELMS] tile-major interleaved fp16
// Processes rows [start_row, end_row) for multi-thread slicing.
// Full range: start_row=0, end_row=n_cols.
static inline void hmx_interleave_rows_to_tiles(__fp16 * restrict vtcm_dst,
const __fp16 * restrict vtcm_src,
int n_cols,
int k,
int src_stride,
int start_row,
int end_row) {
assert(k % HMX_FP16_TILE_N_COLS == 0);
const int n_k_tiles = k / HMX_FP16_TILE_N_COLS;
const HVX_Vector v_scat_base = hvx_vmem(hmx_transpose_scatter_offsets);
const HVX_Vector v_scat_step = Q6_V_vsplat_R(4);
const HVX_VectorPred q_mask64 = Q6_Q_vsetq_R(64);
// Each hvx_vmemu load brings 64 fp16 = 128 bytes covering 2 adjacent K-tiles.
// When n_k_tiles is even, scatter into 2 K-tiles per call (region 4095, no mask)
// using the upper half of hmx_transpose_scatter_offsets. Tail one K-tile (when
// n_k_tiles is odd) falls back to single-tile masked scatter.
const bool pair_scatter = (n_k_tiles & 1) == 0;
const size_t pair_region = (size_t) (2 * HMX_FP16_TILE_SIZE - 1);
const size_t single_region = (size_t) (HMX_FP16_TILE_SIZE - 1);
__builtin_assume(k > 0);
__builtin_assume(end_row > start_row);
if (pair_scatter) {
// Step c by 64 fp16 (two K-tiles per scatter), advance dst by 2 tiles per iter.
const int c_step = 2 * HMX_FP16_TILE_N_COLS;
const size_t c_byte_step = (size_t) c_step * sizeof(__fp16);
const size_t dst_step = 2 * (size_t) HMX_FP16_TILE_N_ELMS;
const int n_c_iters = k / c_step;
for (int r = start_row; r < end_row; r += 2) {
const int ct = r / HMX_FP16_TILE_N_ROWS;
const int local_r = r % HMX_FP16_TILE_N_ROWS;
const bool next_row_valid = (r + 1) < end_row && (r + 1) < n_cols;
const HVX_Vector v_off0 = Q6_Vw_vadd_VwVw(v_scat_base, Q6_V_vsplat_R(local_r * 4));
const HVX_Vector v_off1 = Q6_Vw_vadd_VwVw(v_off0, v_scat_step);
__fp16 * tile_base = vtcm_dst + (size_t) ct * n_k_tiles * HMX_FP16_TILE_N_ELMS;
const uint8_t * p0 = (const uint8_t *) (vtcm_src + r * src_stride);
const uint8_t * p1 = next_row_valid ? (const uint8_t *) (vtcm_src + (r + 1) * src_stride) : NULL;
if (p1) {
for (int i = 0; i < n_c_iters; ++i) {
HVX_Vector v0 = hvx_vmemu(p0);
p0 += c_byte_step;
HVX_Vector v1 = hvx_vmemu(p1);
p1 += c_byte_step;
Q6_vscatter_RMVwV((size_t) tile_base, pair_region, v_off0, v0);
Q6_vscatter_RMVwV((size_t) tile_base, pair_region, v_off1, v1);
tile_base += dst_step;
}
} else {
const HVX_Vector vzero = Q6_V_vzero();
for (int i = 0; i < n_c_iters; ++i) {
HVX_Vector v0 = hvx_vmemu(p0);
p0 += c_byte_step;
Q6_vscatter_RMVwV((size_t) tile_base, pair_region, v_off0, v0);
Q6_vscatter_RMVwV((size_t) tile_base, pair_region, v_off1, vzero);
tile_base += dst_step;
}
}
}
} else {
// Fallback: scatter one K-tile per call (region 2047, masked).
const int c_step = HMX_FP16_TILE_N_COLS;
const size_t c_byte_step = (size_t) c_step * sizeof(__fp16);
const size_t dst_step = (size_t) HMX_FP16_TILE_N_ELMS;
const int n_c_iters = k / c_step;
for (int r = start_row; r < end_row; r += 2) {
const int ct = r / HMX_FP16_TILE_N_ROWS;
const int local_r = r % HMX_FP16_TILE_N_ROWS;
const bool next_row_valid = (r + 1) < end_row && (r + 1) < n_cols;
const HVX_Vector v_off0 = Q6_Vw_vadd_VwVw(v_scat_base, Q6_V_vsplat_R(local_r * 4));
const HVX_Vector v_off1 = Q6_Vw_vadd_VwVw(v_off0, v_scat_step);
__fp16 * tile_base = vtcm_dst + (size_t) ct * n_k_tiles * HMX_FP16_TILE_N_ELMS;
const uint8_t * p0 = (const uint8_t *) (vtcm_src + r * src_stride);
const uint8_t * p1 = next_row_valid ? (const uint8_t *) (vtcm_src + (r + 1) * src_stride) : NULL;
if (p1) {
for (int i = 0; i < n_c_iters; ++i) {
HVX_Vector v0 = hvx_vmemu(p0);
p0 += c_byte_step;
HVX_Vector v1 = hvx_vmemu(p1);
p1 += c_byte_step;
Q6_vscatter_QRMVwV(q_mask64, (size_t) tile_base, single_region, v_off0, v0);
Q6_vscatter_QRMVwV(q_mask64, (size_t) tile_base, single_region, v_off1, v1);
tile_base += dst_step;
}
} else {
const HVX_Vector vzero = Q6_V_vzero();
for (int i = 0; i < n_c_iters; ++i) {
HVX_Vector v0 = hvx_vmemu(p0);
p0 += c_byte_step;
Q6_vscatter_QRMVwV(q_mask64, (size_t) tile_base, single_region, v_off0, v0);
Q6_vscatter_QRMVwV(q_mask64, (size_t) tile_base, single_region, v_off1, vzero);
tile_base += dst_step;
}
}
}
}
}
// Interleave row-major FP16 data into column-major tile format.
// Input: [n_rows, head_dim] row-major. Output: tile[dim_tile][row_tile].
// Processes rows [start_row, end_row) for multi-thread slicing.
// Full range: start_row=0, end_row=n_rows.
static inline void hmx_interleave_cols_to_tiles(__fp16 * restrict tiles_out,
const __fp16 * restrict src,
int n_rows,
int head_dim,
int src_stride,
int n_row_tiles,
int start_row,
int end_row) {
__builtin_assume(head_dim > 0);
const size_t tile_stride_elms = (size_t) n_row_tiles * HMX_FP16_TILE_N_ELMS;
for (int r = start_row; r < end_row; r += 2) {
const bool next_row_valid = (r + 1) < end_row && (r + 1) < n_rows;
const HVX_Vector * pv_in0 = (const HVX_Vector *) (src + r * src_stride);
const HVX_Vector * pv_in1 = next_row_valid ? (const HVX_Vector *) (src + (r + 1) * src_stride) : NULL;
// Row-pair invariants hoisted out of the c loop.
const int r0 = r / HMX_FP16_TILE_N_ROWS;
const int r1_half = (r % HMX_FP16_TILE_N_ROWS) / 2;
// tb0 starts at tile (c0=0, r0); tb1 at the adjacent dim-tile (c0=1, r0).
// Each c step (+= 64) advances both by 2 dim-tiles worth of fp16.
__fp16 * tb0 = tiles_out + (size_t) r0 * HMX_FP16_TILE_N_ELMS;
__fp16 * tb1 = tb0 + tile_stride_elms;
const size_t tb_step = 2 * tile_stride_elms;
if (pv_in1) {
for (int c = 0; c < head_dim; c += 64) {
HVX_Vector v0 = *pv_in0++;
HVX_Vector v1 = *pv_in1++;
HVX_VectorPair vp = Q6_W_vshuff_VVR(v1, v0, -2);
((HVX_Vector *) tb0)[r1_half] = Q6_V_lo_W(vp);
((HVX_Vector *) tb1)[r1_half] = Q6_V_hi_W(vp);
tb0 += tb_step;
tb1 += tb_step;
}
} else {
const HVX_Vector vzero = Q6_V_vzero();
for (int c = 0; c < head_dim; c += 64) {
HVX_Vector v0 = *pv_in0++;
HVX_VectorPair vp = Q6_W_vshuff_VVR(vzero, v0, -2);
((HVX_Vector *) tb0)[r1_half] = Q6_V_lo_W(vp);
((HVX_Vector *) tb1)[r1_half] = Q6_V_hi_W(vp);
tb0 += tb_step;
tb1 += tb_step;
}
}
}
}
#endif // HMX_UTILS_H
+6
View File
@@ -77,6 +77,12 @@ static inline int32_t hvx_vec_get_i32(HVX_Vector v) {
return x;
}
static inline _Float16 hvx_vec_get_f16(HVX_Vector v) {
_Float16 __attribute__((aligned(128))) x;
hvx_vec_store_a(&x, 2, v);
return x;
}
static inline HVX_Vector hvx_vec_abs_f16(HVX_Vector v) {
// abs by clearing the fp16 sign bit
HVX_Vector mask = Q6_Vh_vsplat_R(0x7fff);
+27 -10
View File
@@ -7,7 +7,8 @@
#include "hvx-base.h"
#define hvx_splat_loop_body(dst_type, vec_store) \
#define hvx_splat_pragma(x) _Pragma(#x)
#define hvx_splat_loop_body(dst_type, vec_store, unroll_cnt) \
do { \
dst_type * restrict vdst = (dst_type *) dst; \
\
@@ -16,7 +17,7 @@
\
uint32_t i = 0; \
\
_Pragma("unroll(4)") \
hvx_splat_pragma(unroll(unroll_cnt)) \
for (; i < nvec; i++) { \
vdst[i] = src; \
} \
@@ -25,31 +26,47 @@
} \
} while(0)
static inline void hvx_splat_a(uint8_t * restrict dst, HVX_Vector src, uint32_t n, uint32_t elem_size) {
static inline void hvx_splat_a(void * restrict dst, HVX_Vector src, uint32_t n, uint32_t elem_size) {
assert((unsigned long) dst % 128 == 0);
hvx_splat_loop_body(HVX_Vector, hvx_vec_store_a);
hvx_splat_loop_body(HVX_Vector, hvx_vec_store_a, 4);
}
static inline void hvx_splat_u(uint8_t * restrict dst, HVX_Vector src, uint32_t n, uint32_t elem_size) {
hvx_splat_loop_body(HVX_UVector, hvx_vec_store_u);
static inline void hvx_splat_u(void * restrict dst, HVX_Vector src, uint32_t n, uint32_t elem_size) {
hvx_splat_loop_body(HVX_UVector, hvx_vec_store_u, 4);
}
static inline void hvx_splat_f32_a(uint8_t * restrict dst, float v, uint32_t n) {
static inline void hvx_splat_f32_a(void * restrict dst, float v, uint32_t n) {
hvx_splat_a(dst, hvx_vec_splat_f32(v), n, sizeof(float));
}
static inline void hvx_splat_f32_u(uint8_t * restrict dst, float v, uint32_t n) {
static inline void hvx_splat_f32_u(void * restrict dst, float v, uint32_t n) {
hvx_splat_u(dst, hvx_vec_splat_f32(v), n, sizeof(float));
}
static inline void hvx_splat_f16_a(uint8_t * restrict dst, _Float16 v, uint32_t n) {
static inline void hvx_splat_f16_a(void * restrict dst, _Float16 v, uint32_t n) {
hvx_splat_u(dst, hvx_vec_splat_f16(v), n, sizeof(__fp16));
}
static inline void hvx_splat_f16_u(uint8_t * restrict dst, _Float16 v, uint32_t n) {
static inline void hvx_splat_f16_u(void * restrict dst, _Float16 v, uint32_t n) {
hvx_splat_u(dst, hvx_vec_splat_f16(v), n, sizeof(__fp16));
}
static inline void hvx_splat_u16_a(void * restrict dst, uint16_t v, uint32_t n) {
hvx_splat_a(dst, Q6_Vh_vsplat_R(v), n, sizeof(uint16_t));
}
static inline void hvx_splat_u16_u(void * restrict dst, uint16_t v, uint32_t n) {
hvx_splat_u(dst, Q6_Vh_vsplat_R(v), n, sizeof(uint16_t));
}
static inline void hvx_splat_u8_a(void * restrict dst, uint8_t v, uint32_t n) {
hvx_splat_a(dst, Q6_Vb_vsplat_R(v), n, 1);
}
static inline void hvx_splat_u8_u(void * restrict dst, uint8_t v, uint32_t n) {
hvx_splat_u(dst, Q6_Vb_vsplat_R(v), n, 1);
}
#define hvx_copy_loop_body(dst_type, src_type, vec_store) \
do { \
dst_type * restrict vdst = (dst_type *) dst; \
+2 -2
View File
@@ -17,7 +17,7 @@
#define EXP_LOGN2 (0x3F317218) // ln(2) = 0.6931471805
#define EXP_LOG2E (0x3FB8AA3B) // log2(e) = 1/ln(2) = 1.4426950408
#define EXP_ONE (0x3f800000) // 1.0
#define EXP_RANGE_R (0x42B16666) // 88.7
#define EXP_RANGE_R (0x42B17218) // ln(FLT_MAX) approx = 88.7228
#define EXP_RANGE_L (0xC2B00000) // -88.0 (approx log(FLT_MIN))
static inline HVX_Vector hvx_vec_exp_f32(HVX_Vector in_vec) {
@@ -163,7 +163,7 @@ static inline void hvx_exp_f32(uint8_t * restrict dst, const uint8_t * restrict
HVX_Vector vec_out = Q6_V_vzero();
static const float kInf = INFINITY;
static const float kMaxExp = 88.7f;
static const float kMaxExp = 88.7228f;
const HVX_Vector max_exp = hvx_vec_splat_f32(kMaxExp);
const HVX_Vector inf = hvx_vec_splat_f32(kInf);
+81 -29
View File
@@ -26,8 +26,8 @@ struct htp_unary_context {
const uint8_t * data_src0;
uint8_t * data_dst;
size_t src0_row_size;
size_t dst_row_size;
size_t src0_data_row_size; // actual data bytes per row
size_t dst_data_row_size; // actual data bytes per row
size_t src0_row_size_aligned;
size_t dst_row_size_aligned;
@@ -41,6 +41,40 @@ struct htp_unary_context {
uint32_t nc;
};
// Convert flat row index to DDR byte offset using the tensor's actual strides.
// ir = i1 + ne1*(i2 + ne2*i3) => offset = i1*nb1 + i2*nb2 + i3*nb3
static inline size_t unary_row_offset(uint32_t ir,
uint32_t ne1, uint32_t ne2,
size_t nb1, size_t nb2, size_t nb3) {
const uint32_t i1 = ir % ne1;
const uint32_t i2 = (ir / ne1) % ne2;
const uint32_t i3 = ir / (ne1 * ne2);
return i1 * nb1 + i2 * nb2 + i3 * nb3;
}
// Safe DMA block size from row `ir`: clamp to the tighter dim-1 slice
// boundary of src and dst so the nb1 stride stays valid for all rows.
static inline uint32_t unary_block_size(uint32_t ir,
uint32_t end_row,
uint32_t block,
bool src_contig,
bool dst_contig,
uint32_t src_ne1,
uint32_t dst_ne1) {
uint32_t limit = MIN(block, end_row - ir);
if (!src_contig) {
const uint32_t src_slice_end = (ir / src_ne1 + 1) * src_ne1;
limit = MIN(limit, src_slice_end - ir);
}
if (!dst_contig) {
const uint32_t dst_slice_end = (ir / dst_ne1 + 1) * dst_ne1;
limit = MIN(limit, dst_slice_end - ir);
}
return limit;
}
#define htp_unary_preamble \
const uint32_t ne00 = src->ne[0]; \
const uint32_t ne01 = src->ne[1]; \
@@ -276,8 +310,8 @@ static void unary_job_f32_per_thread(unsigned int nth, unsigned int ith, void *
int32_t * op_params = octx->op_params;
uint32_t src0_nrows_per_thread = uctx->src0_nrows_per_thread;
const size_t src0_row_size = uctx->src0_row_size;
const size_t dst_row_size = uctx->dst_row_size;
const size_t src0_data_row_size = uctx->src0_data_row_size;
const size_t dst_data_row_size = uctx->dst_data_row_size;
const size_t src0_row_size_aligned = uctx->src0_row_size_aligned;
const size_t dst_row_size_aligned = uctx->dst_row_size_aligned;
@@ -303,7 +337,16 @@ static void unary_job_f32_per_thread(unsigned int nth, unsigned int ith, void *
size_t src0_spad_half_size = uctx->src0_spad_half_size;
size_t dst_spad_half_size = uctx->dst_spad_half_size;
const int BLOCK = uctx->block;
// Non-contiguous tensors have gaps at dim-2/3 boundaries that a single-stride
// 2D DMA descriptor cannot span. Clamp BLOCK to ne1 (one dim-1 slice) so every
// transfer stays within a nb1-uniform region. Skipped for contiguous tensors.
const bool src0_contig = (nb02 == (size_t)ne01 * nb01) &&
(nb03 == (size_t)ne02 * nb02);
const bool dst_contig = (nb2 == (size_t)ne1 * nb1) &&
(nb3 == (size_t)ne2 * nb2);
const uint32_t src0_max_block = src0_contig ? uctx->block : MIN((uint32_t)uctx->block, ne01);
const uint32_t dst_max_block = dst_contig ? uctx->block : MIN((uint32_t)uctx->block, ne1);
const uint32_t BLOCK = MIN(src0_max_block, dst_max_block);
if (BLOCK == 0) {
FARF(ERROR, "unary-f32 : current VTCM reservation %zu is too small for even 1 row per thread, needed at least %zu\n",
octx->src0_spad.size_per_thread, src0_row_size_aligned);
@@ -312,21 +355,23 @@ static void unary_job_f32_per_thread(unsigned int nth, unsigned int ith, void *
dma_queue * dma_queue = octx->ctx->dma[ith];
for (uint32_t ir = src0_start_row, spad_idx = 0; ir < src0_end_row && spad_idx < 2; ir += BLOCK, spad_idx++) {
const uint32_t block_size = MIN(BLOCK, src0_end_row - ir);
for (uint32_t ir = src0_start_row, spad_idx = 0; ir < src0_end_row && spad_idx < 2; spad_idx++) {
const uint32_t block_size = unary_block_size(ir, src0_end_row, BLOCK, src0_contig, dst_contig, ne01, ne1);
// Dummy DMA transation for sequencing (interleaving dst,src,dst,...)
dma_queue_push_vtcm_to_ddr(dma_queue,
dma_queue_push(dma_queue,
dma_make_ptr(data_dst, dst_spad_data + (spad_idx * dst_spad_half_size)),
dst_row_size, dst_row_size_aligned, 0);
nb1, dst_row_size_aligned, dst_data_row_size, 0);
dma_queue_push_ddr_to_vtcm(dma_queue,
dma_make_ptr(src0_spad_data + (spad_idx * src0_spad_half_size), data_src + (ir * src0_row_size)),
src0_row_size_aligned, src0_row_size, block_size);
const size_t src0_off = unary_row_offset(ir, ne01, ne02, nb01, nb02, nb03);
dma_queue_push(dma_queue,
dma_make_ptr(src0_spad_data + (spad_idx * src0_spad_half_size), data_src + src0_off),
src0_row_size_aligned, nb01, src0_data_row_size, block_size);
ir += block_size;
}
for (uint32_t ir = src0_start_row; ir < src0_end_row; ir += BLOCK) {
const uint32_t block_size = MIN(BLOCK, src0_end_row - ir);
for (uint32_t ir = src0_start_row; ir < src0_end_row; ) {
const uint32_t block_size = unary_block_size(ir, src0_end_row, BLOCK, src0_contig, dst_contig, ne01, ne1);
float * dst_spad = (float *) dma_queue_pop(dma_queue).src;
float * src0_spad = (float *) dma_queue_pop(dma_queue).dst;
@@ -361,18 +406,25 @@ static void unary_job_f32_per_thread(unsigned int nth, unsigned int ith, void *
break;
}
dma_queue_push_vtcm_to_ddr(dma_queue,
dma_make_ptr(data_dst + (ir * dst_row_size), dst_spad),
dst_row_size, dst_row_size_aligned, block_size);
const size_t dst_off = unary_row_offset(ir, ne1, ne2, nb1, nb2, nb3);
dma_queue_push(dma_queue,
dma_make_ptr(data_dst + dst_off, dst_spad),
nb1, dst_row_size_aligned, dst_data_row_size, block_size);
// prefetch N+2 loop iteration if any
const uint32_t pref_block = (ir + BLOCK * 2);
if (pref_block < src0_end_row) {
const uint32_t pref_block_size = MIN(BLOCK, src0_end_row - pref_block);
dma_queue_push_ddr_to_vtcm(dma_queue,
dma_make_ptr(src0_spad, data_src + (pref_block * src0_row_size)),
src0_row_size_aligned, src0_row_size, pref_block_size);
const uint32_t next_ir = ir + block_size;
if (next_ir < src0_end_row) {
const uint32_t next_block_size = unary_block_size(next_ir, src0_end_row, BLOCK, src0_contig, dst_contig, ne01, ne1);
const uint32_t pref_ir = next_ir + next_block_size;
if (pref_ir < src0_end_row) {
const uint32_t pref_block_size = unary_block_size(pref_ir, src0_end_row, BLOCK, src0_contig, dst_contig, ne01, ne1);
const size_t src0_pref_off = unary_row_offset(pref_ir, ne01, ne02, nb01, nb02, nb03);
dma_queue_push(dma_queue,
dma_make_ptr(src0_spad, data_src + src0_pref_off),
src0_row_size_aligned, nb01, src0_data_row_size, pref_block_size);
}
}
ir += block_size;
}
dma_queue_flush(dma_queue);
@@ -426,11 +478,11 @@ static int execute_op_unary_f32(struct htp_ops_context * octx) {
const uint32_t src0_nrows = src0->ne[1] * src0->ne[2] * src0->ne[3];
const uint32_t n_threads = MIN(octx->n_threads, src0_nrows);
const size_t src0_row_size = src0->nb[1];
const size_t dst_row_size = dst->nb[1];
const size_t src0_data_row_size = src0->ne[0] * sizeof(float);
const size_t dst_data_row_size = dst->ne[0] * sizeof(float);
const size_t src0_row_size_aligned = hex_round_up(src0_row_size, VLEN);
const size_t dst_row_size_aligned = hex_round_up(dst_row_size, VLEN);
const size_t src0_row_size_aligned = hex_round_up(src0_data_row_size, VLEN);
const size_t dst_row_size_aligned = hex_round_up(dst_data_row_size, VLEN);
// VTCM scratchpads for all tensors
// N rows per thread, padded to HVX vector size
@@ -468,8 +520,8 @@ static int execute_op_unary_f32(struct htp_ops_context * octx) {
.data_src0 = (const uint8_t *)src0->data,
.data_dst = (uint8_t *)dst->data,
.src0_row_size = src0_row_size,
.dst_row_size = dst_row_size,
.src0_data_row_size = src0_data_row_size,
.dst_data_row_size = dst_data_row_size,
.src0_row_size_aligned = src0_row_size_aligned,
.dst_row_size_aligned = dst_row_size_aligned,
+16
View File
@@ -0,0 +1,16 @@
#ifndef VTCM_UTILS_H
#define VTCM_UTILS_H
#include "hex-utils.h"
#include <assert.h>
#include <stdint.h>
#include <hexagon_types.h>
static inline uint8_t *vtcm_seq_alloc(uint8_t **vtcm_ptr, size_t size) {
uint8_t *p = *vtcm_ptr;
*vtcm_ptr += size;
return p;
}
#endif // VTCM_UTILS_H
+3 -3
View File
@@ -166,8 +166,8 @@ static ggml_backend_buffer_i ggml_backend_metal_buffer_private_i = {
/* .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,
/* .get_tensor_2d_async = */ NULL,
/* .set_tensor_2d_async = */ NULL,
/* .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,
@@ -567,8 +567,8 @@ static ggml_backend_i ggml_backend_metal_i = {
/* .free = */ ggml_backend_metal_free,
/* .set_tensor_async = */ ggml_backend_metal_set_tensor_async,
/* .get_tensor_async = */ ggml_backend_metal_get_tensor_async,
/* .get_tensor_2d_async = */ NULL,
/* .set_tensor_2d_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .cpy_tensor_async = */ ggml_backend_metal_cpy_tensor_async, // only needed for multi-GPU setups
/* .synchronize = */ ggml_backend_metal_synchronize,
/* .graph_plan_create = */ NULL,
+2 -2
View File
@@ -4343,9 +4343,9 @@ static ggml_backend_i ggml_backend_opencl_i = {
/* .free = */ ggml_backend_opencl_free,
/* .set_tensor_async = */ NULL, /* ggml_backend_opencl_set_tensor_async */
/* .get_tensor_async = */ NULL, /* ggml_backend_opencl_get_tensor_async */
/* .cpy_tensor_async = */ NULL, /* ggml_backend_opencl_cpy_tensor_async */
/* .get_tensor_2d_async = */ NULL,
/* .set_tensor_2d_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .cpy_tensor_async = */ NULL, /* ggml_backend_opencl_cpy_tensor_async */
/* .synchronize = */ ggml_backend_opencl_synchronize,
/* .graph_plan_create = */ NULL,
/* .graph_plan_free = */ NULL,
+2 -2
View File
@@ -740,9 +740,9 @@ static ggml_backend_i ggml_backend_rpc_interface = {
/* .free = */ ggml_backend_rpc_free,
/* .set_tensor_async = */ NULL,
/* .get_tensor_async = */ NULL,
/* .cpy_tensor_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .set_tensor_2d_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .cpy_tensor_async = */ NULL,
/* .synchronize = */ ggml_backend_rpc_synchronize,
/* .graph_plan_create = */ NULL,
/* .graph_plan_free = */ NULL,
+1 -1
View File
@@ -4700,8 +4700,8 @@ static ggml_backend_i ggml_backend_sycl_interface = {
/* .free = */ ggml_backend_sycl_free,
/* .set_tensor_async = */ ggml_backend_sycl_set_tensor_async,
/* .get_tensor_async = */ ggml_backend_sycl_get_tensor_async,
/* .get_tensor_2d_async = */ NULL,
/* .set_tensor_2d_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .cpy_tensor_async = */ NULL, // ggml_backend_sycl_cpy_tensor_async,
// // TODO: update for the new
// interface
+1 -1
View File
@@ -34,8 +34,8 @@ static ggml_backend_i ggml_backend_remoting_interface = {
/* .free = */ ggml_backend_remoting_free,
/* .set_tensor_async = */ NULL, // ggml_backend_remoting_set_tensor_async,
/* .get_tensor_async = */ NULL, // ggml_backend_remoting_get_tensor_async,
/* .get_tensor_2d_async = */ NULL,
/* .set_tensor_2d_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .cpy_tensor_async = */ NULL, // ggml_backend_remoting_cpy_tensor_async,
/* .synchronize = */ NULL, // ggml_backend_remoting_synchronize,
/* .graph_plan_create = */ NULL,
+261 -109
View File
@@ -440,10 +440,12 @@ struct vk_fa_pipeline_state {
bool f32acc;
uint32_t flags;
uint32_t limit_occupancy_shmem;
ggml_type k_type;
ggml_type v_type;
bool operator<(const vk_fa_pipeline_state &b) const {
return std::tie(HSK, HSV, Br, Bc, D_split, row_split, shmem_staging, path, workgroup_size, subgroup_size, aligned, f32acc, flags, limit_occupancy_shmem) <
std::tie(b.HSK, b.HSV, b.Br, b.Bc, b.D_split, b.row_split, b.shmem_staging, b.path, b.workgroup_size, b.subgroup_size, b.aligned, b.f32acc, b.flags, b.limit_occupancy_shmem);
return std::tie(HSK, HSV, Br, Bc, D_split, row_split, shmem_staging, path, workgroup_size, subgroup_size, aligned, f32acc, flags, limit_occupancy_shmem, k_type, v_type) <
std::tie(b.HSK, b.HSV, b.Br, b.Bc, b.D_split, b.row_split, b.shmem_staging, b.path, b.workgroup_size, b.subgroup_size, b.aligned, b.f32acc, b.flags, b.limit_occupancy_shmem, b.k_type, b.v_type);
}
};
@@ -3041,7 +3043,7 @@ static vk_fa_tuning_params get_fa_tuning_params_coopmat1(const vk_device& device
return result;
}
static vk_fa_tuning_params get_fa_tuning_params_coopmat2(const vk_device& device, uint32_t hsk, uint32_t hsv, uint32_t n_rows, uint32_t n_kv, ggml_type kv_type, bool f32acc) {
static vk_fa_tuning_params get_fa_tuning_params_coopmat2(const vk_device& device, uint32_t hsk, uint32_t hsv, uint32_t n_rows, uint32_t n_kv, ggml_type k_type, ggml_type v_type, bool f32acc) {
GGML_UNUSED(n_kv);
GGML_UNUSED(f32acc);
@@ -3055,7 +3057,7 @@ static vk_fa_tuning_params get_fa_tuning_params_coopmat2(const vk_device& device
if (small_rows) {
result.block_rows = 32;
result.block_cols = 32;
} else if (ggml_is_quantized(kv_type) || hsk >= 256 || hsv >= 256) {
} else if (ggml_is_quantized(k_type) || ggml_is_quantized(v_type) || hsk >= 256 || hsv >= 256) {
result.block_rows = (hsk >= 512 || hsv >= 512) ? 32 : 64;
result.block_cols = 32;
} else {
@@ -3069,7 +3071,13 @@ static vk_fa_tuning_params get_fa_tuning_params_coopmat2(const vk_device& device
return result;
}
static vk_fa_tuning_params get_fa_tuning_params(const vk_device& device, uint32_t hsk, uint32_t hsv, uint32_t n_rows, uint32_t n_kv, ggml_type kv_type, bool f32acc) {
static vk_fa_tuning_params get_fa_tuning_params(const vk_device& device, uint32_t hsk, uint32_t hsv, uint32_t n_rows, uint32_t n_kv, ggml_type k_type, ggml_type v_type, bool f32acc) {
// Mixed K/V is only implemented on the coopmat2 (flash_attn_cm2) path; never use scalar/cm1.
if (k_type != v_type) {
GGML_ASSERT(device->coopmat2);
return get_fa_tuning_params_coopmat2(device, hsk, hsv, n_rows, n_kv, k_type, v_type, f32acc);
}
FaCodePath path = device->coopmat2 ? FA_COOPMAT2 :
device->coopmat1_fa_support ? FA_COOPMAT1 : FA_SCALAR;
@@ -3081,7 +3089,7 @@ static vk_fa_tuning_params get_fa_tuning_params(const vk_device& device, uint32_
if (path == FA_COOPMAT1) {
bool shape_ok = (f32acc && device->coopmat_support_16x16x16_f32acc) ||
(!f32acc && device->coopmat_support_16x16x16_f16acc);
const vk_fa_tuning_params params = get_fa_tuning_params_coopmat1(device, hsk, hsv, n_rows, n_kv, kv_type, f32acc);
const vk_fa_tuning_params params = get_fa_tuning_params_coopmat1(device, hsk, hsv, n_rows, n_kv, k_type, f32acc);
bool shmem_ok = ggml_vk_flash_attn_coopmat_shmem_support(device, params, hsk, hsv, f32acc);
if (!shape_ok || !shmem_ok) {
@@ -3094,20 +3102,25 @@ static vk_fa_tuning_params get_fa_tuning_params(const vk_device& device, uint32_
path = FA_SCALAR;
}
// Q1_0 K/V is only implemented on coopmat2 (flash_attn_cm2); there is no scalar FA shader for it.
if ((k_type == GGML_TYPE_Q1_0 || v_type == GGML_TYPE_Q1_0) && device->coopmat2) {
path = FA_COOPMAT2;
}
switch (path) {
case FA_SCALAR:
return get_fa_tuning_params_scalar(device, hsk, hsv, n_rows, n_kv, kv_type, f32acc);
return get_fa_tuning_params_scalar(device, hsk, hsv, n_rows, n_kv, k_type, f32acc);
case FA_COOPMAT1:
return get_fa_tuning_params_coopmat1(device, hsk, hsv, n_rows, n_kv, kv_type, f32acc);
return get_fa_tuning_params_coopmat1(device, hsk, hsv, n_rows, n_kv, k_type, f32acc);
case FA_COOPMAT2:
return get_fa_tuning_params_coopmat2(device, hsk, hsv, n_rows, n_kv, kv_type, f32acc);
return get_fa_tuning_params_coopmat2(device, hsk, hsv, n_rows, n_kv, k_type, v_type, f32acc);
default:
throw std::runtime_error("unsupported FaCodePath");
}
}
static vk_fa_pipeline_state get_fa_pipeline_state(const vk_device& device, const vk_fa_tuning_params& params, uint32_t hsk, uint32_t hsv, bool aligned, bool f32acc,
bool use_mask, bool use_mask_opt, bool use_logit_softcap) {
bool use_mask, bool use_mask_opt, bool use_logit_softcap, ggml_type k_type, ggml_type v_type) {
const bool old_amd_windows = device->vendor_id == VK_VENDOR_ID_AMD && device->driver_id == vk::DriverId::eAmdProprietary &&
(device->architecture == AMD_GCN || device->architecture == AMD_RDNA1 || device->architecture == AMD_RDNA2);
@@ -3118,12 +3131,32 @@ static vk_fa_pipeline_state get_fa_pipeline_state(const vk_device& device, const
const uint32_t subgroup_size = params.disable_subgroups ? 0 : params.subgroup_size;
return vk_fa_pipeline_state{hsk, hsv, params.block_rows, params.block_cols, params.d_split, params.row_split, params.shmem_staging, params.path, params.workgroup_size, subgroup_size, aligned, f32acc, flags, params.limit_occupancy_shmem};
return vk_fa_pipeline_state{hsk, hsv, params.block_rows, params.block_cols, params.d_split, params.row_split, params.shmem_staging, params.path, params.workgroup_size, subgroup_size, aligned, f32acc, flags, params.limit_occupancy_shmem, k_type, v_type};
}
static std::vector<uint32_t> get_fa_spec_constants(const vk_fa_pipeline_state& state) {
return {state.workgroup_size, state.Br, state.Bc, state.HSK, state.HSV, !state.aligned, state.D_split,
state.row_split, state.subgroup_size, state.shmem_staging ? 1u : 0u, state.flags, state.limit_occupancy_shmem};
const auto fa_block_bytes = [](ggml_type t) -> uint32_t {
// decodeBufF32 uses a block of vec4s for a better memory access pattern.
return t == GGML_TYPE_F32 ? 16u : (uint32_t) ggml_type_size(t);
};
return {
/* 0 WorkGroupSize */ state.workgroup_size,
/* 1 Br */ state.Br,
/* 2 Bc */ state.Bc,
/* 3 HSK */ state.HSK,
/* 4 HSV */ state.HSV,
/* 5 Clamp */ static_cast<uint32_t>(!state.aligned),
/* 6 D_split */ state.D_split,
/* 7 row_split */ state.row_split,
/* 8 SubGroupSize */ state.subgroup_size,
/* 9 SHMEM_STAGING */ state.shmem_staging ? 1u : 0u,
/*10 Flags */ state.flags,
/*11 LIMIT_OCCUPANCY_SHMEM */ state.limit_occupancy_shmem,
/*12 FaTypeK */ static_cast<uint32_t>(state.k_type),
/*13 FaTypeV */ static_cast<uint32_t>(state.v_type),
/*14 FaBlockBytesK */ fa_block_bytes(state.k_type),
/*15 FaBlockBytesV */ fa_block_bytes(state.v_type),
};
}
static bool ggml_vk_matmul_shmem_support(const vk_device& device, const std::vector<uint32_t>& warptile, bool mul_mat_id, ggml_type src0_type) {
@@ -3578,16 +3611,35 @@ static void ggml_vk_load_shaders(vk_device& device) {
}
#endif
#if defined(VK_NV_cooperative_matrix2) && defined(GGML_VULKAN_COOPMAT2_GLSLC_SUPPORT)
#define CREATE_FA_CM2_MIXED() \
for (int fa_k_ty = 0; fa_k_ty < (int)GGML_TYPE_COUNT; ++fa_k_ty) { \
for (auto &fa : device->pipeline_flash_attn_f32_f16[fa_k_ty]) { \
FaCodePath path = fa.first.path; \
uint32_t Br = fa.first.Br; \
uint32_t Bc = fa.first.Bc; \
bool aligned = fa.first.aligned; \
bool f32acc = fa.first.f32acc; \
if (path == FA_COOPMAT2) { \
if (aligned) { \
if (f32acc) { \
ggml_vk_create_pipeline(device, fa.second, "flash_attn_f32_f16_mixed_aligned_f32acc_cm2", flash_attn_f32_f16_mixed_cm2_len, flash_attn_f32_f16_mixed_cm2_data, "main", 7, sizeof(vk_flash_attn_push_constants), {Br, 1, 1}, get_fa_spec_constants(fa.first), Bc, true, false, 0); \
} else { \
ggml_vk_create_pipeline(device, fa.second, "flash_attn_f32_f16_mixed_aligned_f16acc_cm2", flash_attn_f32_f16_mixed_f16acc_cm2_len, flash_attn_f32_f16_mixed_f16acc_cm2_data, "main", 7, sizeof(vk_flash_attn_push_constants), {Br, 1, 1}, get_fa_spec_constants(fa.first), Bc, true, false, 0); \
} \
} else { \
if (f32acc) { \
ggml_vk_create_pipeline(device, fa.second, "flash_attn_f32_f16_mixed_f32acc_cm2", flash_attn_f32_f16_mixed_cm2_len, flash_attn_f32_f16_mixed_cm2_data, "main", 7, sizeof(vk_flash_attn_push_constants), {Br, 1, 1}, get_fa_spec_constants(fa.first), 1, true, false, 0); \
} else { \
ggml_vk_create_pipeline(device, fa.second, "flash_attn_f32_f16_mixed_f16acc_cm2", flash_attn_f32_f16_mixed_f16acc_cm2_len, flash_attn_f32_f16_mixed_f16acc_cm2_data, "main", 7, sizeof(vk_flash_attn_push_constants), {Br, 1, 1}, get_fa_spec_constants(fa.first), 1, true, false, 0); \
} \
} \
} \
} \
}
if (device->coopmat2) {
CREATE_FA(GGML_TYPE_F32, f32, FA_COOPMAT2, _cm2)
CREATE_FA(GGML_TYPE_F16, f16, FA_COOPMAT2, _cm2)
CREATE_FA(GGML_TYPE_Q4_0, q4_0, FA_COOPMAT2, _cm2)
CREATE_FA(GGML_TYPE_Q4_1, q4_1, FA_COOPMAT2, _cm2)
CREATE_FA(GGML_TYPE_Q5_0, q5_0, FA_COOPMAT2, _cm2)
CREATE_FA(GGML_TYPE_Q5_1, q5_1, FA_COOPMAT2, _cm2)
CREATE_FA(GGML_TYPE_Q8_0, q8_0, FA_COOPMAT2, _cm2)
CREATE_FA(GGML_TYPE_IQ4_NL, iq4_nl, FA_COOPMAT2, _cm2)
CREATE_FA_CM2_MIXED();
}
#undef CREATE_FA_CM2_MIXED
#endif
#undef CREATE_FA
@@ -6845,7 +6897,7 @@ static void ggml_vk_buffer_write_nc_async(ggml_backend_vk_context * ctx, vk_cont
}
}
static bool ggml_vk_buffer_write_2d_async(vk_context subctx, vk_buffer& dst, size_t offset, const void * src, size_t spitch, size_t width, size_t height, bool sync_staging = false) {
static bool ggml_vk_buffer_write_2d_async(vk_context subctx, vk_buffer& dst, size_t offset, const void * src, size_t spitch, size_t dpitch, size_t width, size_t height, bool sync_staging = false) {
VK_LOG_DEBUG("ggml_vk_buffer_write_2d_async(" << width << ", " << height << ")");
// Check if src is pinned memory
vk_buffer buf = nullptr;
@@ -6855,7 +6907,7 @@ static bool ggml_vk_buffer_write_2d_async(vk_context subctx, vk_buffer& dst, siz
if (buf != nullptr) {
// Memory is pinned, use as staging buffer
std::vector<vk::BufferCopy> slices(1);
if (width == spitch) {
if (width == spitch && width == dpitch) {
// Only do single write if stride is equal
slices[0].srcOffset = buf_offset;
slices[0].dstOffset = offset;
@@ -6864,7 +6916,7 @@ static bool ggml_vk_buffer_write_2d_async(vk_context subctx, vk_buffer& dst, siz
slices.resize(height);
for (size_t i = 0; i < height; i++) {
slices[i].srcOffset = buf_offset + i * spitch;
slices[i].dstOffset = offset + i * width;
slices[i].dstOffset = offset + i * dpitch;
slices[i].size = width;
}
}
@@ -6881,21 +6933,30 @@ static bool ggml_vk_buffer_write_2d_async(vk_context subctx, vk_buffer& dst, siz
}
// Staging buffer required
const size_t copy_size = width*height;
ggml_vk_ensure_sync_staging_buffer(dst->device, copy_size);
const size_t staging_size = width * height;
ggml_vk_ensure_sync_staging_buffer(dst->device, staging_size);
vk_buffer& staging_buffer = dst->device->sync_staging;
VkBufferCopy buf_copy = {
0,
offset,
copy_size};
std::vector<vk::BufferCopy> slices(1);
if (width == dpitch) {
slices[0].srcOffset = 0;
slices[0].dstOffset = offset;
slices[0].size = staging_size;
} else {
slices.resize(height);
for (size_t i = 0; i < height; i++) {
slices[i].srcOffset = i * width;
slices[i].dstOffset = offset + i * dpitch;
slices[i].size = width;
}
}
ggml_vk_sync_buffers(nullptr, subctx);
vkCmdCopyBuffer(subctx->s->buffer->buf, (VkBuffer)staging_buffer->buffer, (VkBuffer)dst->buffer, 1, &buf_copy);
subctx->s->buffer->buf.copyBuffer((VkBuffer)staging_buffer->buffer, (VkBuffer)dst->buffer, slices);
if (width == spitch) {
deferred_memcpy((uint8_t *)staging_buffer->ptr, src, width * height, &subctx->in_memcpys);
deferred_memcpy((uint8_t *)staging_buffer->ptr, src, staging_size, &subctx->in_memcpys);
} else {
for (size_t i = 0; i < height; i++) {
deferred_memcpy((uint8_t *)staging_buffer->ptr + i * width, (const uint8_t *) src + i * spitch, width, &subctx->in_memcpys);
@@ -6906,24 +6967,24 @@ static bool ggml_vk_buffer_write_2d_async(vk_context subctx, vk_buffer& dst, siz
static bool ggml_vk_buffer_write_async(vk_context subctx, vk_buffer& dst, size_t offset, const void * src, size_t size, bool sync_staging = false) {
VK_LOG_DEBUG("ggml_vk_buffer_write_async(" << size << ")");
return ggml_vk_buffer_write_2d_async(subctx, dst, offset, src, size, size, 1, sync_staging);
return ggml_vk_buffer_write_2d_async(subctx, dst, offset, src, size, size, size, 1, sync_staging);
}
static void ggml_vk_buffer_write_2d(vk_buffer& dst, size_t offset, const void * src, size_t spitch, size_t width, size_t height) {
static void ggml_vk_buffer_write_2d(vk_buffer& dst, size_t offset, const void * src, size_t spitch, size_t dpitch, size_t width, size_t height) {
VK_LOG_DEBUG("ggml_vk_buffer_write_2d(" << width << ", " << height << ")");
// Buffer is already mapped
if(dst->memory_property_flags & vk::MemoryPropertyFlagBits::eHostVisible) {
GGML_ASSERT(dst->memory_property_flags & vk::MemoryPropertyFlagBits::eHostCoherent);
for (size_t i = 0; i < height; i++) {
memcpy((uint8_t *)dst->ptr + offset + i * width, (const uint8_t *) src + i * spitch, width);
memcpy((uint8_t *)dst->ptr + offset + i * dpitch, (const uint8_t *) src + i * spitch, width);
}
} else {
std::lock_guard<std::recursive_mutex> guard(dst->device->mutex);
vk_context subctx = ggml_vk_create_temporary_context(dst->device->transfer_queue.cmd_pool);
ggml_vk_ctx_begin(dst->device, subctx);
bool ret = ggml_vk_buffer_write_2d_async(subctx, dst, offset, src, spitch, width, height, true);
bool ret = ggml_vk_buffer_write_2d_async(subctx, dst, offset, src, spitch, dpitch, width, height, true);
GGML_ASSERT(ret);
ggml_vk_ctx_end(subctx);
@@ -6944,7 +7005,7 @@ static void ggml_vk_buffer_write_2d(vk_buffer& dst, size_t offset, const void *
static void ggml_vk_buffer_write(vk_buffer& dst, size_t offset, const void * src, size_t size) {
VK_LOG_DEBUG("ggml_vk_buffer_write(" << size << ")");
ggml_vk_buffer_write_2d(dst, offset, src, 0, size, 1);
ggml_vk_buffer_write_2d(dst, offset, src, size, size, size, 1);
}
static bool ggml_vk_buffer_read_2d_async(vk_context subctx, vk_buffer& src, size_t offset, void * dst, size_t spitch, size_t dpitch, size_t width, size_t height, bool sync_staging = false) {
@@ -6990,15 +7051,35 @@ static bool ggml_vk_buffer_read_2d_async(vk_context subctx, vk_buffer& src, size
}
// Fall back to staging buffer
const size_t copy_size = dpitch * height;
ggml_vk_ensure_sync_staging_buffer(src->device, copy_size);
const size_t staging_size = width * height;
ggml_vk_ensure_sync_staging_buffer(src->device, staging_size);
vk_buffer& staging_buffer = src->device->sync_staging;
ggml_vk_sync_buffers(nullptr, subctx);
subctx->s->buffer->buf.copyBuffer(src->buffer, staging_buffer->buffer, slices);
std::vector<vk::BufferCopy> staging_slices(1);
if (width == spitch) {
staging_slices[0].srcOffset = offset;
staging_slices[0].dstOffset = 0;
staging_slices[0].size = staging_size;
} else {
staging_slices.resize(height);
for (size_t i = 0; i < height; i++) {
staging_slices[i].srcOffset = offset + i * spitch;
staging_slices[i].dstOffset = i * width;
staging_slices[i].size = width;
}
}
deferred_memcpy(dst, staging_buffer->ptr, copy_size, &subctx->out_memcpys);
ggml_vk_sync_buffers(nullptr, subctx);
subctx->s->buffer->buf.copyBuffer(src->buffer, staging_buffer->buffer, staging_slices);
if (width == dpitch) {
deferred_memcpy(dst, staging_buffer->ptr, staging_size, &subctx->out_memcpys);
} else {
for (size_t i = 0; i < height; i++) {
deferred_memcpy((uint8_t *) dst + i * dpitch, (const uint8_t *) staging_buffer->ptr + i * width, width, &subctx->out_memcpys);
}
}
return true;
}
@@ -7006,8 +7087,8 @@ static bool ggml_vk_buffer_read_async(vk_context subctx, vk_buffer& src, size_t
return ggml_vk_buffer_read_2d_async(subctx, src, offset, dst, size, size, size, 1, sync_staging);
}
static void ggml_vk_buffer_read(vk_buffer& src, size_t offset, void * dst, size_t size) {
VK_LOG_DEBUG("ggml_vk_buffer_read(" << src->buffer << ", " << offset << ", " << size << ")");
static void ggml_vk_buffer_read_2d(vk_buffer& src, size_t offset, void * dst, size_t spitch, size_t dpitch, size_t width, size_t height) {
VK_LOG_DEBUG("ggml_vk_buffer_read_2d(" << src->buffer << ", " << offset << ", " << width << ", " << height << ")");
// If the device is not an UMA device the memory is host-accessible through rebar. While writing
// through PCIe is sufficient fast reading back data from PCIe is slower than going through
@@ -7015,18 +7096,20 @@ static void ggml_vk_buffer_read(vk_buffer& src, size_t offset, void * dst, size_
if(src->memory_property_flags & vk::MemoryPropertyFlagBits::eHostVisible && src->device->uma) {
GGML_ASSERT(src->memory_property_flags & vk::MemoryPropertyFlagBits::eHostCoherent);
memcpy(dst, (uint8_t *) src->ptr + offset, size);
for (size_t i = 0; i < height; i++) {
memcpy((uint8_t *) dst + i * dpitch, (const uint8_t *) src->ptr + offset + i * spitch, width);
}
} else {
std::lock_guard<std::recursive_mutex> guard(src->device->mutex);
vk_context subctx = ggml_vk_create_temporary_context(src->device->transfer_queue.cmd_pool);
ggml_vk_ctx_begin(src->device, subctx);
bool ret = ggml_vk_buffer_read_async(subctx, src, offset, dst, size, true);
bool ret = ggml_vk_buffer_read_2d_async(subctx, src, offset, dst, spitch, dpitch, width, height, true);
GGML_ASSERT(ret);
ggml_vk_ctx_end(subctx);
ggml_vk_submit(subctx, src->device->fence);
VK_CHECK(src->device->device.waitForFences({ src->device->fence }, true, UINT64_MAX), "vk_buffer_read waitForFences");
VK_CHECK(src->device->device.waitForFences({ src->device->fence }, true, UINT64_MAX), "vk_buffer_read_2d waitForFences");
src->device->device.resetFences({ src->device->fence });
ggml_vk_queue_command_pools_cleanup(src->device);
@@ -7036,6 +7119,11 @@ static void ggml_vk_buffer_read(vk_buffer& src, size_t offset, void * dst, size_
}
}
static void ggml_vk_buffer_read(vk_buffer& src, size_t offset, void * dst, size_t size) {
VK_LOG_DEBUG("ggml_vk_buffer_read(" << src->buffer << ", " << offset << ", " << size << ")");
ggml_vk_buffer_read_2d(src, offset, dst, size, size, size, 1);
}
static void ggml_vk_buffer_copy_async(vk_context& ctx, vk_buffer& dst, size_t dst_offset, vk_buffer& src, size_t src_offset, size_t size) {
VK_LOG_DEBUG("ggml_vk_buffer_copy_async(" << size << ")");
// Make sure both buffers are on same device
@@ -7067,7 +7155,7 @@ static void ggml_vk_buffer_copy(vk_buffer& dst, size_t dst_offset, vk_buffer& sr
// Copy to src staging buffer
ggml_vk_buffer_copy(src->device->sync_staging, 0, src, src_offset, size);
// Copy to dst buffer
ggml_vk_buffer_write_2d(dst, dst_offset, src->device->sync_staging->ptr, 0, size, 1);
ggml_vk_buffer_write(dst, dst_offset, src->device->sync_staging->ptr, size);
}
}
@@ -9006,8 +9094,6 @@ static void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx
assert(dst->type == GGML_TYPE_F32);
assert(q->type == GGML_TYPE_F32);
assert(k->type == v->type);
uint32_t gqa_ratio = 1;
uint32_t qk_ratio = neq2 / nek2;
uint32_t workgroups_x = (uint32_t)neq1;
@@ -9018,7 +9104,7 @@ static void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx
// For scalar/coopmat1 FA, we can use the "large" size to accommodate qga.
// For coopmat2 FA, we always use the small size (which is still pretty large for gqa).
vk_fa_tuning_params tuning_params = get_fa_tuning_params(ctx->device, HSK, HSV, 512, KV, k->type, f32acc);
vk_fa_tuning_params tuning_params = get_fa_tuning_params(ctx->device, HSK, HSV, 512, KV, k->type, v->type, f32acc);
const uint32_t max_gqa = std::min(tuning_params.block_rows, 32u);
if (N <= 8 && qk_ratio > 1 && qk_ratio <= max_gqa &&
@@ -9031,7 +9117,11 @@ static void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx
workgroups_y /= gqa_ratio;
}
tuning_params = get_fa_tuning_params(ctx->device, HSK, HSV, N, KV, k->type, f32acc);
tuning_params = get_fa_tuning_params(ctx->device, HSK, HSV, N, KV, k->type, v->type, f32acc);
if (tuning_params.path != FA_COOPMAT2) {
GGML_ASSERT(k->type == v->type);
}
const uint32_t q_stride = (uint32_t)(nbq1 / ggml_type_size(q->type));
uint32_t k_stride = (uint32_t)(nbk1 / ggml_type_size(k->type));
@@ -9070,7 +9160,7 @@ static void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx
// Only use mask opt when the mask is fairly large. This hasn't been tuned extensively.
bool use_mask_opt = mask && nem1 >= 32 && nem0 * nem1 > 32768 && nem0 >= tuning_params.block_cols * 16;
vk_fa_pipeline_state fa_pipeline_state = get_fa_pipeline_state(ctx->device, tuning_params, HSK, HSV, aligned, f32acc,
mask != nullptr, use_mask_opt, logit_softcap != 0);
mask != nullptr, use_mask_opt, logit_softcap != 0, k->type, v->type);
vk_pipeline pipeline = nullptr;
@@ -13615,6 +13705,20 @@ static void ggml_backend_vk_buffer_set_tensor(ggml_backend_buffer_t buffer, ggml
ggml_vk_buffer_write(buf, vk_tensor_offset(tensor) + tensor->view_offs + offset, data, size);
}
static void ggml_backend_vk_buffer_set_tensor_2d(ggml_backend_buffer_t buffer, ggml_tensor * tensor, const void * data, size_t offset,
size_t size, size_t n_copies, size_t stride_tensor, size_t stride_data) {
VK_LOG_DEBUG("ggml_backend_vk_buffer_set_tensor_2d(" << buffer << ", " << tensor << ", " << data << ", " << offset << ", " << size << ", " <<
n_copies << ", " << stride_tensor << ", " << stride_data << ")");
ggml_backend_vk_buffer_context * buf_ctx = (ggml_backend_vk_buffer_context *)buffer->context;
vk_buffer buf = buf_ctx->dev_buffer;
if (size == 0) {
return;
}
ggml_vk_buffer_write_2d(buf, vk_tensor_offset(tensor) + tensor->view_offs + offset, data, stride_data, stride_tensor, size, n_copies);
}
static void ggml_backend_vk_buffer_get_tensor(ggml_backend_buffer_t buffer, const ggml_tensor * tensor, void * data, size_t offset, size_t size) {
VK_LOG_DEBUG("ggml_backend_vk_buffer_get_tensor(" << buffer << ", " << tensor << ", " << data << ", " << offset << ", " << size << ")");
ggml_backend_vk_buffer_context * buf_ctx = (ggml_backend_vk_buffer_context *)buffer->context;
@@ -13628,6 +13732,21 @@ static void ggml_backend_vk_buffer_get_tensor(ggml_backend_buffer_t buffer, cons
ggml_vk_buffer_read(buf, vk_tensor_offset(tensor) + tensor->view_offs + offset, data, size);
}
static void ggml_backend_vk_buffer_get_tensor_2d(ggml_backend_buffer_t buffer, const ggml_tensor * tensor, void * data, size_t offset,
size_t size, size_t n_copies, size_t stride_tensor, size_t stride_data) {
VK_LOG_DEBUG("ggml_backend_vk_buffer_get_tensor_2d(" << buffer << ", " << tensor << ", " << data << ", " << offset << ", " << size << ", " <<
n_copies << ", " << stride_tensor << ", " << stride_data << ")");
ggml_backend_vk_buffer_context * buf_ctx = (ggml_backend_vk_buffer_context *)buffer->context;
if (size == 0) {
return;
}
vk_buffer buf = buf_ctx->dev_buffer;
ggml_vk_buffer_read_2d(buf, vk_tensor_offset(tensor) + tensor->view_offs + offset, data, stride_tensor, stride_data, size, n_copies);
}
static bool ggml_backend_vk_buffer_cpy_tensor(ggml_backend_buffer_t buffer, const ggml_tensor * src, ggml_tensor * dst) {
if (ggml_nbytes(src) == 0) {
return true;
@@ -13662,8 +13781,8 @@ static ggml_backend_buffer_i ggml_backend_vk_buffer_interface = {
/* .memset_tensor = */ ggml_backend_vk_buffer_memset_tensor,
/* .set_tensor = */ ggml_backend_vk_buffer_set_tensor,
/* .get_tensor = */ ggml_backend_vk_buffer_get_tensor,
/* .set_tensor_2d = */ NULL,
/* .get_tensor_2d = */ NULL,
/* .set_tensor_2d = */ ggml_backend_vk_buffer_set_tensor_2d,
/* .get_tensor_2d = */ ggml_backend_vk_buffer_get_tensor_2d,
/* .cpy_tensor = */ ggml_backend_vk_buffer_cpy_tensor,
/* .clear = */ ggml_backend_vk_buffer_clear,
/* .reset = */ NULL,
@@ -13819,8 +13938,9 @@ static ggml_backend_buffer_type_t ggml_backend_vk_get_default_buffer_type(ggml_b
return &ctx->device->buffer_type;
}
static void ggml_backend_vk_set_tensor_async(ggml_backend_t backend, ggml_tensor * tensor, const void * data, size_t offset, size_t size) {
VK_LOG_DEBUG("ggml_backend_vk_set_tensor_async(" << size << ")");
static void ggml_backend_vk_set_tensor_2d_async(ggml_backend_t backend, ggml_tensor * tensor, const void * data, size_t offset,
size_t size, size_t n_copies, size_t stride_tensor, size_t stride_data) {
VK_LOG_DEBUG("ggml_backend_vk_set_tensor_2d_async(" << size << ", " << n_copies << ")");
ggml_backend_vk_context * ctx = (ggml_backend_vk_context *)backend->context;
GGML_ASSERT((tensor->buffer->buft == ggml_backend_vk_get_default_buffer_type(backend) || tensor->buffer->buft == ggml_backend_vk_host_buffer_type()) && "unsupported buffer type");
@@ -13834,7 +13954,6 @@ static void ggml_backend_vk_set_tensor_async(ggml_backend_t backend, ggml_tensor
if (ctx->device->async_use_transfer_queue) {
if (ctx->transfer_ctx.expired()) {
// Initialize new transfer context
cpy_ctx = ggml_vk_create_context(ctx, ctx->transfer_cmd_pool);
ctx->transfer_ctx = cpy_ctx;
ggml_vk_ctx_begin(ctx->device, cpy_ctx);
@@ -13849,25 +13968,48 @@ static void ggml_backend_vk_set_tensor_async(ggml_backend_t backend, ggml_tensor
auto dst_offset = vk_tensor_offset(tensor) + tensor->view_offs + offset;
bool ret = ggml_vk_buffer_write_async(cpy_ctx, buf, dst_offset, data, size);
bool ret = ggml_vk_buffer_write_2d_async(cpy_ctx, buf, dst_offset, data, stride_data, stride_tensor, size, n_copies);
if (!ret) {
ggml_vk_ensure_sync_staging_buffer(ctx, size);
const size_t staging_size = size * n_copies;
ggml_vk_ensure_sync_staging_buffer(ctx, staging_size);
ggml_vk_sync_buffers(nullptr, cpy_ctx);
vk::BufferCopy buffer_cpy;
buffer_cpy.srcOffset = 0;
buffer_cpy.dstOffset = dst_offset;
buffer_cpy.size = size;
std::vector<vk::BufferCopy> slices(1);
if (size == stride_tensor) {
slices[0].srcOffset = 0;
slices[0].dstOffset = dst_offset;
slices[0].size = staging_size;
} else {
slices.resize(n_copies);
for (size_t i = 0; i < n_copies; i++) {
slices[i].srcOffset = i * size;
slices[i].dstOffset = dst_offset + i * stride_tensor;
slices[i].size = size;
}
}
cpy_ctx->s->buffer->buf.copyBuffer(ctx->sync_staging->buffer, buf->buffer, { buffer_cpy });
deferred_memcpy(ctx->sync_staging->ptr, data, size, &cpy_ctx->in_memcpys);
cpy_ctx->s->buffer->buf.copyBuffer(ctx->sync_staging->buffer, buf->buffer, slices);
if (size == stride_data) {
deferred_memcpy(ctx->sync_staging->ptr, data, staging_size, &cpy_ctx->in_memcpys);
} else {
for (size_t i = 0; i < n_copies; i++) {
deferred_memcpy((uint8_t *)ctx->sync_staging->ptr + i * size, (const uint8_t *)data + i * stride_data, size, &cpy_ctx->in_memcpys);
}
}
ggml_vk_synchronize(ctx);
}
}
static void ggml_backend_vk_get_tensor_async(ggml_backend_t backend, const ggml_tensor * tensor, void * data, size_t offset, size_t size) {
VK_LOG_DEBUG("ggml_backend_vk_get_tensor_async(" << size << ")");
static void ggml_backend_vk_set_tensor_async(ggml_backend_t backend, ggml_tensor * tensor, const void * data, size_t offset, size_t size) {
VK_LOG_DEBUG("ggml_backend_vk_set_tensor_async(" << size << ")");
ggml_backend_vk_set_tensor_2d_async(backend, tensor, data, offset, size, 1, size, size);
}
static void ggml_backend_vk_get_tensor_2d_async(ggml_backend_t backend, const ggml_tensor * tensor, void * data, size_t offset,
size_t size, size_t n_copies, size_t stride_tensor, size_t stride_data) {
VK_LOG_DEBUG("ggml_backend_vk_get_tensor_2d_async(" << size << ", " << n_copies << ")");
ggml_backend_vk_context * ctx = (ggml_backend_vk_context *)backend->context;
GGML_ASSERT((tensor->buffer->buft == ggml_backend_vk_get_default_buffer_type(backend) || tensor->buffer->buft == ggml_backend_vk_host_buffer_type()) && "unsupported buffer type");
@@ -13882,24 +14024,45 @@ static void ggml_backend_vk_get_tensor_async(ggml_backend_t backend, const ggml_
vk_buffer buf = buf_ctx->dev_buffer;
auto src_offset = vk_tensor_offset(tensor) + tensor->view_offs + offset;
bool ret = ggml_vk_buffer_read_async(compute_ctx, buf, src_offset, data, size);
bool ret = ggml_vk_buffer_read_2d_async(compute_ctx, buf, src_offset, data, stride_tensor, stride_data, size, n_copies);
// If that failed, copy synchronously through a staging buffer
if (!ret) {
ggml_vk_ensure_sync_staging_buffer(ctx, size);
const size_t staging_size = size * n_copies;
ggml_vk_ensure_sync_staging_buffer(ctx, staging_size);
ggml_vk_sync_buffers(nullptr, compute_ctx);
vk::BufferCopy buffer_cpy;
buffer_cpy.srcOffset = src_offset;
buffer_cpy.dstOffset = 0;
buffer_cpy.size = size;
std::vector<vk::BufferCopy> slices(1);
if (size == stride_tensor) {
slices[0].srcOffset = src_offset;
slices[0].dstOffset = 0;
slices[0].size = staging_size;
} else {
slices.resize(n_copies);
for (size_t i = 0; i < n_copies; i++) {
slices[i].srcOffset = src_offset + i * stride_tensor;
slices[i].dstOffset = i * size;
slices[i].size = size;
}
}
compute_ctx->s->buffer->buf.copyBuffer(buf->buffer, ctx->sync_staging->buffer, { buffer_cpy });
deferred_memcpy(data, ctx->sync_staging->ptr, size, &compute_ctx->out_memcpys);
compute_ctx->s->buffer->buf.copyBuffer(buf->buffer, ctx->sync_staging->buffer, slices);
if (size == stride_data) {
deferred_memcpy(data, ctx->sync_staging->ptr, staging_size, &compute_ctx->out_memcpys);
} else {
for (size_t i = 0; i < n_copies; i++) {
deferred_memcpy((uint8_t *)data + i * stride_data, (const uint8_t *)ctx->sync_staging->ptr + i * size, size, &compute_ctx->out_memcpys);
}
}
ggml_vk_synchronize(ctx);
}
}
static void ggml_backend_vk_get_tensor_async(ggml_backend_t backend, const ggml_tensor * tensor, void * data, size_t offset, size_t size) {
VK_LOG_DEBUG("ggml_backend_vk_get_tensor_async(" << size << ")");
ggml_backend_vk_get_tensor_2d_async(backend, tensor, data, offset, size, 1, size, size);
}
static bool ggml_backend_vk_cpy_tensor_async(ggml_backend_t backend_src, ggml_backend_t backend_dst, const ggml_tensor * src, ggml_tensor * dst) {
VK_LOG_DEBUG("ggml_backend_vk_cpy_tensor_async(" << src << " -> " << dst << ", size=" << ggml_nbytes(src) << ")");
ggml_backend_vk_context * ctx = (ggml_backend_vk_context *)backend_dst->context;
@@ -15123,8 +15286,8 @@ static ggml_backend_i ggml_backend_vk_interface = {
/* .free = */ ggml_backend_vk_free,
/* .set_tensor_async = */ ggml_backend_vk_set_tensor_async,
/* .get_tensor_async = */ ggml_backend_vk_get_tensor_async,
/* .get_tensor_2d_async = */ NULL,
/* .set_tensor_2d_async = */ NULL,
/* .set_tensor_2d_async = */ ggml_backend_vk_set_tensor_2d_async,
/* .get_tensor_2d_async = */ ggml_backend_vk_get_tensor_2d_async,
/* .cpy_tensor_async = */ ggml_backend_vk_cpy_tensor_async,
/* .synchronize = */ ggml_backend_vk_synchronize,
/* .graph_plan_create = */ NULL,
@@ -15481,38 +15644,27 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
if (op->src[3] && op->src[3]->type != GGML_TYPE_F16) {
return false;
}
// It's straightforward to support different K/V dequant, but would
// significantly increase the number of pipelines
if (op->src[1]->type != op->src[2]->type) {
// mismatching K/V type is currently supported for coopmat2 only.
if (op->src[1]->type != op->src[2]->type && !coopmat2) {
return false;
}
switch (op->src[1]->type) {
case GGML_TYPE_F16:
case GGML_TYPE_F32:
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q8_0:
case GGML_TYPE_Q4_1:
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1:
case GGML_TYPE_IQ4_NL:
// supported in scalar and coopmat2 paths
break;
// K dequants currently disabled because D dimension is rounded up to 256 and runs inefficiently
//case GGML_TYPE_Q2_K:
//case GGML_TYPE_Q3_K:
//case GGML_TYPE_Q4_K:
//case GGML_TYPE_Q5_K:
//case GGML_TYPE_Q6_K:
//case GGML_TYPE_IQ1_S:
//case GGML_TYPE_IQ1_M:
//case GGML_TYPE_IQ2_XXS:
//case GGML_TYPE_IQ2_XS:
//case GGML_TYPE_IQ2_S:
//case GGML_TYPE_IQ3_XXS:
//case GGML_TYPE_IQ3_S:
//case GGML_TYPE_IQ4_XS:
default:
auto fa_kv_ok = [coopmat2](ggml_type t) {
switch (t) {
case GGML_TYPE_F32:
case GGML_TYPE_F16:
case GGML_TYPE_Q8_0:
case GGML_TYPE_Q5_1:
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q4_1:
case GGML_TYPE_Q4_0:
return true;
case GGML_TYPE_Q1_0:
return coopmat2;
default:
return false;
}
};
if (!fa_kv_ok(op->src[1]->type) || !fa_kv_ok(op->src[2]->type)) {
return false;
}
if (!coopmat2 && !(device->subgroup_shuffle && device->subgroup_vote)) {
@@ -13,6 +13,12 @@ layout (constant_id = 8) const uint32_t SubGroupSize = 32;
layout (constant_id = 9) const uint32_t SHMEM_STAGING = 0;
layout (constant_id = 10) const uint32_t Flags = 0;
layout (constant_id = 11) const uint32_t LIMIT_OCCUPANCY_SHMEM = 0;
// ggml_type enumerant for K/V
layout (constant_id = 12) const uint32_t FaTypeK = 0;
layout (constant_id = 13) const uint32_t FaTypeV = 0;
// sizeof(decode buffer): quants -> ggml block size; F32 -> 16 (decodeBufF32 vec4).
layout (constant_id = 14) const uint32_t FaBlockBytesK = 2;
layout (constant_id = 15) const uint32_t FaBlockBytesV = 2;
const bool USE_MASK_OPT = (Flags & 1) != 0;
const bool MASK_ENABLE = (Flags & 2) != 0;
@@ -17,8 +17,57 @@
#extension GL_EXT_null_initializer : enable
#include "types.glsl"
#include "dequant_funcs_cm2.glsl"
#include "flash_attn_base.glsl"
#include "dequant_funcs_cm2.glsl"
// buffer_reference stride = sizeof(struct) = FaBlockBytesK/V.
layout(buffer_reference, std430, buffer_reference_align = 1) buffer decodeBufFA_K {
uint8_t raw[FaBlockBytesK];
};
layout(buffer_reference, std430, buffer_reference_align = 1) buffer decodeBufFA_V {
uint8_t raw[FaBlockBytesV];
};
uint fa_block_elems(uint ty) {
switch (ty) {
case 0u: return 4u; // GGML_TYPE_F32: vec4 block (matches decodeBufF32 / dequantFuncF32)
case 1u: return 1u; // GGML_TYPE_F16
case 2u: return uint(QUANT_K_Q4_0);
case 3u: return uint(QUANT_K_Q4_1);
case 6u: return uint(QUANT_K_Q5_0);
case 7u: return uint(QUANT_K_Q5_1);
case 8u: return uint(QUANT_K_Q8_0);
case 41u: return uint(QUANT_K_Q1_0);
default:
return 1u;
}
}
float16_t faDecodeK(const decodeBufFA_K bl_in, const uint blockCoords[2], const uint coordInBlock[2]) {
switch (FaTypeK) {
case 0u: return dequantFuncF32(decodeBufF32(bl_in), blockCoords, coordInBlock);
case 2u: return dequantFuncQ4_0(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
case 3u: return dequantFuncQ4_1(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
case 6u: return dequantFuncQ5_0(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
case 7u: return dequantFuncQ5_1(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
case 8u: return dequantFuncQ8_0(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
case 41u: return dequantFuncQ1_0(decodeBufQ1_0(bl_in), blockCoords, coordInBlock);
default: return float16_t(0);
}
}
float16_t faDecodeV(const decodeBufFA_V bl_in, const uint blockCoords[2], const uint coordInBlock[2]) {
switch (FaTypeV) {
case 0u: return dequantFuncF32(decodeBufF32(bl_in), blockCoords, coordInBlock);
case 2u: return dequantFuncQ4_0(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
case 3u: return dequantFuncQ4_1(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
case 6u: return dequantFuncQ5_0(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
case 7u: return dequantFuncQ5_1(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
case 8u: return dequantFuncQ8_0(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
case 41u: return dequantFuncQ1_0(decodeBufQ1_0(bl_in), blockCoords, coordInBlock);
default: return float16_t(0);
}
}
layout (binding = 0) readonly buffer Q {uint8_t data_q[];};
layout (binding = 1) readonly buffer K {uint8_t data_k[];};
@@ -55,12 +104,6 @@ ACC_TYPE Max(const in uint32_t row, const in uint32_t col, const in ACC_TYPE ele
return max(elem0, elem1);
}
#if BLOCK_SIZE > 1
#define DECODEFUNC , DEQUANTFUNC
#else
#define DECODEFUNC
#endif
// Store the output when doing grouped query attention.
// Rows index by Q's dimension 2, and the first N rows are valid.
D_TYPE perElemOpGqaStore(const in uint32_t r, const in uint32_t c, const in D_TYPE elem, const in uint32_t o_offset, const in uint32_t iq2, const in uint32_t N)
@@ -95,10 +138,6 @@ ACC_TYPE perElemOpNonGqaSplitKStoreCol0(const in uint32_t r, const in uint32_t c
}
void main() {
#ifdef NEEDS_INIT_IQ_SHMEM
init_iq_shmem(gl_WorkGroupSize);
#endif
init_indices();
tensorLayoutNV<2, gl_CooperativeMatrixClampModeConstantNV> tensorLayoutQ = createTensorLayoutNV(2, gl_CooperativeMatrixClampModeConstantNV);
@@ -107,10 +146,10 @@ void main() {
tensorViewNV<2, false, 1, 0> tensorViewTranspose = createTensorViewNV(2, false, 1, 0);
#if BLOCK_SIZE > 1
tensorLayoutK = setTensorLayoutBlockSizeNV(tensorLayoutK, 1, BLOCK_SIZE);
tensorLayoutV = setTensorLayoutBlockSizeNV(tensorLayoutV, 1, BLOCK_SIZE);
#endif
const uint bs_k = fa_block_elems(FaTypeK);
const uint bs_v = fa_block_elems(FaTypeV);
tensorLayoutK = setTensorLayoutBlockSizeNV(tensorLayoutK, 1, bs_k);
tensorLayoutV = setTensorLayoutBlockSizeNV(tensorLayoutV, 1, bs_v);
tensorLayoutQ = setTensorLayoutDimensionNV(tensorLayoutQ, N, HSK);
tensorLayoutK = setTensorLayoutDimensionNV(tensorLayoutK, KV, HSK);
@@ -120,10 +159,12 @@ void main() {
if (Clamp != gl_CooperativeMatrixClampModeConstantNV)
{
q_stride &= ~7;
#if BLOCK_SIZE == 1
k_stride &= ~7;
v_stride &= ~7;
#endif
if (bs_k == 1u) {
k_stride &= ~7;
}
if (bs_v == 1u) {
v_stride &= ~7;
}
m_stride &= ~7;
}
tensorLayoutQ = setTensorLayoutStrideNV(tensorLayoutQ, q_stride, 1);
@@ -230,7 +271,13 @@ void main() {
coopmat<float16_t, gl_ScopeWorkgroup, HSK_pad, Bc, gl_MatrixUseB> K_T;
uint32_t k_offset = ik2*p.nb12 + ik3*p.nb13;
coopMatLoadTensorNV(K_T, data_k, k_offset, sliceTensorLayoutNV(tensorLayoutK, j * Bc, Bc, 0, HSK_pad), tensorViewTranspose DECODEFUNC);
// F16: bs_k==1 (direct load). F32: bs_k==4 (vec4 / dequantFuncF32). Q4/Q8 family: bs_k==32. Q1_0: bs_k==128.
const bool k_use_decode = (bs_k > 1u);
if (k_use_decode) {
coopMatLoadTensorNV(K_T, data_k, k_offset, sliceTensorLayoutNV(tensorLayoutK, j * Bc, Bc, 0, HSK_pad), tensorViewTranspose, faDecodeK);
} else {
coopMatLoadTensorNV(K_T, data_k, k_offset, sliceTensorLayoutNV(tensorLayoutK, j * Bc, Bc, 0, HSK_pad), tensorViewTranspose);
}
S = coopMatMulAdd(Qf16, K_T, S);
if (LOGIT_SOFTCAP) {
@@ -291,7 +338,12 @@ void main() {
coopmat<float16_t, gl_ScopeWorkgroup, Bc, HSV_pad, gl_MatrixUseB> V;
uint32_t v_offset = iv2*p.nb22 + iv3*p.nb23;
coopMatLoadTensorNV(V, data_v, v_offset, sliceTensorLayoutNV(tensorLayoutV, j * Bc, Bc, 0, HSV_pad) DECODEFUNC);
const bool v_use_decode = (bs_v > 1u);
if (v_use_decode) {
coopMatLoadTensorNV(V, data_v, v_offset, sliceTensorLayoutNV(tensorLayoutV, j * Bc, Bc, 0, HSV_pad), faDecodeV);
} else {
coopMatLoadTensorNV(V, data_v, v_offset, sliceTensorLayoutNV(tensorLayoutV, j * Bc, Bc, 0, HSV_pad));
}
L = eM*L + rowsum;
@@ -641,20 +641,17 @@ void process_shaders() {
fa_base_dict["ACC_TYPE_MAX"] = "float16_t(65504.0)";
}
if (fp16) {
#if defined(GGML_VULKAN_COOPMAT2_GLSLC_SUPPORT)
string_to_spv("flash_attn_f32_f16_mixed", "flash_attn_cm2.comp",
merge_maps(fa_base_dict, {{"Q_TYPE", "float"}, {"D_TYPE", "float"}, {"D_TYPEV4", "vec4"}}), fp16, false, true, f16acc);
#endif
}
for (const auto& tname : type_names) {
if (tname == "bf16") continue;
if (fp16) {
#if defined(GGML_VULKAN_COOPMAT2_GLSLC_SUPPORT)
if (tname == "f16") {
string_to_spv("flash_attn_f32_f16_" + tname, "flash_attn_cm2.comp",
merge_maps(fa_base_dict, {{"Q_TYPE", "float"}, {"D_TYPE", "float"}, {"D_TYPEV4", "vec4"}}), fp16, false, true, f16acc);
} else {
std::string data_a_key = "DATA_A_" + to_uppercase(tname);
string_to_spv("flash_attn_f32_f16_" + tname, "flash_attn_cm2.comp",
merge_maps(fa_base_dict, {{data_a_key, "1"}, {"Q_TYPE", "float"}, {"D_TYPE", "float"}, {"D_TYPEV4", "vec4"}, {"DEQUANTFUNC", "dequantFunc"+to_uppercase(tname) }, {"BLOCK_SIZE", "QUANT_K_"+to_uppercase(tname) }}), fp16, false, true, f16acc);
}
#endif
#if defined(GGML_VULKAN_COOPMAT_GLSLC_SUPPORT)
if (tname == "f16") {
string_to_spv("flash_attn_f32_f16_" + tname, "flash_attn_cm1.comp",
+263 -19
View File
@@ -1,6 +1,7 @@
#ifndef GGML_WEBGPU_SHADER_LIB_HPP
#define GGML_WEBGPU_SHADER_LIB_HPP
#include "ggml-impl.h"
#include "ggml-wgsl-shaders.hpp"
#include "ggml.h"
#include "pre_wgsl.hpp"
@@ -405,6 +406,31 @@ struct ggml_webgpu_scale_pipeline_key_hash {
}
};
/** Upscale **/
struct ggml_webgpu_upscale_pipeline_key {
ggml_type input_type;
ggml_type output_type;
uint32_t base_mode;
bool antialias;
bool operator==(const ggml_webgpu_upscale_pipeline_key & other) const {
return input_type == other.input_type && output_type == other.output_type && base_mode == other.base_mode &&
antialias == other.antialias;
}
};
struct ggml_webgpu_upscale_pipeline_key_hash {
size_t operator()(const ggml_webgpu_upscale_pipeline_key & key) const {
size_t seed = 0;
ggml_webgpu_hash_combine(seed, key.input_type);
ggml_webgpu_hash_combine(seed, key.output_type);
ggml_webgpu_hash_combine(seed, key.base_mode);
ggml_webgpu_hash_combine(seed, key.antialias);
return seed;
}
};
/** Concat **/
struct ggml_webgpu_concat_pipeline_key {
@@ -664,7 +690,7 @@ inline uint32_t ggml_webgpu_flash_attn_max_kv_tile(const ggml_webgpu_shader_lib_
}
const size_t base_q_bytes = (key.head_dim_qk + key.head_dim_v) * q_tile * GGML_WEBGPU_F16_SIZE_BYTES +
2 * q_tile * GGML_WEBGPU_F32_SIZE_BYTES;
size_t bytes_per_kv = 0;
size_t bytes_per_kv = 0;
if (!key.kv_direct) {
bytes_per_kv += std::max(key.head_dim_qk, key.head_dim_v);
}
@@ -701,10 +727,10 @@ inline ggml_webgpu_flash_attn_decisions ggml_webgpu_flash_attn_get_decisions(
(v_offset_elems % GGML_WEBGPU_FLASH_ATTN_TILE_KV_VEC_WIDTH == 0u);
const bool kv_vec_type_supported =
K->type == GGML_TYPE_F16 || K->type == GGML_TYPE_Q4_0 || K->type == GGML_TYPE_Q8_0;
const bool use_vec = context.supports_subgroups && (context.src0->ne[1] < 20) && (context.src0->ne[0] % 32 == 0) &&
(context.src2->ne[0] % GGML_WEBGPU_FLASH_ATTN_TILE_KV_VEC_WIDTH == 0) &&
kv_vec_type_supported && (K->type != GGML_TYPE_F16 || f16_vec4_aligned) &&
(context.src2->type == K->type);
const bool use_vec = context.supports_subgroups && (context.src0->ne[1] < 20) && (context.src0->ne[0] % 32 == 0) &&
(context.src2->ne[0] % GGML_WEBGPU_FLASH_ATTN_TILE_KV_VEC_WIDTH == 0) &&
kv_vec_type_supported && (K->type != GGML_TYPE_F16 || f16_vec4_aligned) &&
(context.src2->type == K->type);
const bool use_tile = context.supports_subgroups && !context.supports_subgroup_matrix && K->type == GGML_TYPE_F16 &&
V->type == GGML_TYPE_F16 && f16_vec4_aligned &&
(context.src0->ne[0] % GGML_WEBGPU_FLASH_ATTN_TILE_KV_VEC_WIDTH == 0) &&
@@ -862,9 +888,12 @@ struct ggml_webgpu_mul_mat_shader_decisions {
struct ggml_webgpu_mul_mat_id_pipeline_key {
ggml_type src0_type;
ggml_type src1_type;
uint32_t n_experts;
int vectorized;
bool operator==(const ggml_webgpu_mul_mat_id_pipeline_key & other) const {
return src0_type == other.src0_type && src1_type == other.src1_type;
return src0_type == other.src0_type && src1_type == other.src1_type && n_experts == other.n_experts &&
vectorized == other.vectorized;
}
};
@@ -873,6 +902,8 @@ struct ggml_webgpu_mul_mat_id_pipeline_key_hash {
size_t seed = 0;
ggml_webgpu_hash_combine(seed, key.src0_type);
ggml_webgpu_hash_combine(seed, key.src1_type);
ggml_webgpu_hash_combine(seed, key.n_experts);
ggml_webgpu_hash_combine(seed, key.vectorized);
return seed;
}
};
@@ -1023,6 +1054,8 @@ class ggml_webgpu_shader_lib {
std::unordered_map<int, webgpu_pipeline> mul_mat_id_gather_pipelines; // key is fixed
std::unordered_map<ggml_webgpu_mul_mat_id_pipeline_key, webgpu_pipeline, ggml_webgpu_mul_mat_id_pipeline_key_hash>
mul_mat_id_pipelines; // src0_type/src1_type
std::unordered_map<ggml_webgpu_mul_mat_id_pipeline_key, webgpu_pipeline, ggml_webgpu_mul_mat_id_pipeline_key_hash>
mul_mat_id_vec_pipelines; // src0_type/src1_type
std::unordered_map<ggml_webgpu_set_rows_pipeline_key, webgpu_pipeline, ggml_webgpu_set_rows_pipeline_key_hash>
set_rows_pipelines;
@@ -1042,6 +1075,8 @@ class ggml_webgpu_shader_lib {
webgpu_pipeline,
ggml_webgpu_rms_norm_mul_pipeline_key_hash>
rms_norm_mul_pipelines;
std::unordered_map<ggml_webgpu_upscale_pipeline_key, webgpu_pipeline, ggml_webgpu_upscale_pipeline_key_hash>
upscale_pipelines;
public:
ggml_webgpu_shader_lib(wgpu::Device device) { this->device = device; }
@@ -1516,7 +1551,7 @@ class ggml_webgpu_shader_lib {
key.type = context.dst->type;
key.d_state = (int) context.src0->ne[0];
key.xbc_overlap = ggml_webgpu_tensor_overlap(context.src1, context.src4) &&
ggml_webgpu_tensor_overlap(context.src1, context.src5);
ggml_webgpu_tensor_overlap(context.src1, context.src5);
auto it = ssm_scan_pipelines.find(key);
if (it != ssm_scan_pipelines.end()) {
@@ -1633,10 +1668,10 @@ class ggml_webgpu_shader_lib {
ggml_webgpu_mul_mat_vec_pipeline_key key = {};
key.src0_type = context.src0->type;
key.src1_type = context.src1->type;
key.vectorized = (context.src0->ne[0] % 4 == 0 &&
key.vectorized = (context.src0->ne[0] % 4 == 0 &&
(context.src0->type == GGML_TYPE_F32 || context.src0->type == GGML_TYPE_F16)) ?
1 :
0;
1 :
0;
auto it = mul_mat_vec_pipelines.find(key);
if (it != mul_mat_vec_pipelines.end()) {
@@ -1744,12 +1779,12 @@ class ggml_webgpu_shader_lib {
webgpu_pipeline get_mul_mat_fast_pipeline(const ggml_webgpu_shader_lib_context & context) {
ggml_webgpu_mul_mat_pipeline_key key = {};
key.src0_type = context.src0->type;
key.src1_type = context.src1->type;
key.vectorized = (context.src0->ne[0] % 4 == 0 && context.dst->ne[0] % 4 == 0 && context.dst->ne[1] % 4 == 0 &&
(context.src0->type == GGML_TYPE_F32 || context.src0->type == GGML_TYPE_F16)) ?
1 :
0;
key.src0_type = context.src0->type;
key.src1_type = context.src1->type;
key.vectorized = (context.src0->ne[0] % 4 == 0 && context.dst->ne[0] % 4 == 0 &&
(context.src0->type == GGML_TYPE_F32 || context.src0->type == GGML_TYPE_F16)) ?
1 :
0;
key.use_subgroup_matrix = context.supports_subgroup_matrix;
auto it = mul_mat_fast_pipelines.find(key);
@@ -2012,6 +2047,11 @@ class ggml_webgpu_shader_lib {
ggml_webgpu_mul_mat_id_pipeline_key key = {};
key.src0_type = context.src0->type;
key.src1_type = context.src1->type;
key.n_experts = context.src0->ne[2];
key.vectorized = (context.src0->ne[0] % 4 == 0 && context.src0->ne[1] % 4 == 0 &&
(context.src0->type == GGML_TYPE_F32 || context.src0->type == GGML_TYPE_F16)) ?
1 :
0;
auto it = mul_mat_id_pipelines.find(key);
if (it != mul_mat_id_pipelines.end()) {
@@ -2041,14 +2081,12 @@ class ggml_webgpu_shader_lib {
switch (context.src0->type) {
case GGML_TYPE_F32:
defines.push_back("SRC0_INNER_TYPE=f32");
defines.push_back("FLOAT");
defines.push_back("INIT_SRC0_SHMEM_FLOAT");
defines.push_back("INIT_SRC1_SHMEM_FLOAT");
variant += "_f32";
break;
case GGML_TYPE_F16:
defines.push_back("SRC0_INNER_TYPE=f16");
defines.push_back("FLOAT");
defines.push_back("INIT_SRC0_SHMEM_FLOAT");
defines.push_back("INIT_SRC1_SHMEM_FLOAT");
variant += "_f16";
@@ -2064,12 +2102,32 @@ class ggml_webgpu_shader_lib {
defines.push_back("U32_DEQUANT_HELPERS");
defines.push_back("SRC0_INNER_TYPE=u32");
switch (context.src0->type) {
case GGML_TYPE_IQ1_S:
case GGML_TYPE_IQ1_M:
case GGML_TYPE_IQ4_NL:
case GGML_TYPE_IQ4_XS:
defines.push_back(type_upper + "_GRID");
break;
case GGML_TYPE_IQ2_XXS:
case GGML_TYPE_IQ2_XS:
case GGML_TYPE_IQ2_S:
case GGML_TYPE_IQ3_XXS:
case GGML_TYPE_IQ3_S:
defines.push_back(type_upper + "_GRID");
defines.push_back(type_upper + "_TABLES");
break;
default:
break;
}
variant += std::string("_") + src0_name;
break;
}
}
defines.push_back("SCALAR");
// VEC/SCALAR controls
defines.push_back(key.vectorized ? "VEC" : "SCALAR");
// mul_mat_id is register-tile only.
const uint32_t tile_k =
@@ -2085,6 +2143,9 @@ class ggml_webgpu_shader_lib {
// variant suffix for src1 type
variant += std::string("_") + (context.src1->type == GGML_TYPE_F32 ? "f32" : "f16");
if (key.vectorized) {
variant += "_vectorized";
}
auto processed = preprocessor.preprocess(wgsl_mul_mat_id, defines);
@@ -2102,6 +2163,123 @@ class ggml_webgpu_shader_lib {
return mul_mat_id_pipelines[key];
}
webgpu_pipeline get_mul_mat_id_vec_pipeline(const ggml_webgpu_shader_lib_context & context) {
ggml_webgpu_mul_mat_id_pipeline_key key = {};
key.src0_type = context.src0->type;
key.src1_type = context.src1->type;
key.n_experts = context.src0->ne[2];
key.vectorized = (context.src0->ne[0] % 4 == 0 &&
(context.src0->type == GGML_TYPE_F32 || context.src0->type == GGML_TYPE_F16)) ?
1 :
0;
auto it = mul_mat_id_vec_pipelines.find(key);
if (it != mul_mat_id_vec_pipelines.end()) {
return it->second;
}
std::vector<std::string> defines;
std::string variant = "mul_mat_id_vec";
const char * shader_src = wgsl_mul_mat_id_vec;
// src1 type
switch (context.src1->type) {
case GGML_TYPE_F32:
defines.push_back("SRC1_INNER_TYPE=f32");
break;
case GGML_TYPE_F16:
defines.push_back("SRC1_INNER_TYPE=f16");
break;
default:
GGML_ABORT("Unsupported src1 type for mul_mat fast shader");
}
// src0 type
switch (context.src0->type) {
case GGML_TYPE_F32:
defines.push_back("SRC0_INNER_TYPE=f32");
defines.push_back("MUL_ACC_FLOAT");
variant += "_f32";
break;
case GGML_TYPE_F16:
defines.push_back("SRC0_INNER_TYPE=f16");
defines.push_back("MUL_ACC_FLOAT");
variant += "_f16";
break;
default:
{
// Quantized types: use helpers but accumulate in f16
const struct ggml_type_traits * src0_traits = ggml_get_type_traits(context.src0->type);
std::string src0_name = src0_traits->type_name;
std::string type_upper = src0_name;
variant += "_" + src0_name;
std::transform(type_upper.begin(), type_upper.end(), type_upper.begin(), ::toupper);
defines.push_back("BYTE_HELPERS");
defines.push_back("MUL_ACC_" + type_upper);
defines.push_back("U32_DEQUANT_HELPERS");
defines.push_back("SRC0_INNER_TYPE=u32");
switch (context.src0->type) {
case GGML_TYPE_IQ1_S:
case GGML_TYPE_IQ1_M:
case GGML_TYPE_IQ2_S:
case GGML_TYPE_IQ3_S:
case GGML_TYPE_IQ4_NL:
case GGML_TYPE_IQ4_XS:
defines.push_back(type_upper + "_GRID");
break;
case GGML_TYPE_IQ2_XXS:
case GGML_TYPE_IQ2_XS:
case GGML_TYPE_IQ3_XXS:
defines.push_back(type_upper + "_GRID");
defines.push_back(type_upper + "_TABLES");
break;
default:
break;
}
break;
}
}
// VEC/SCALAR controls
defines.push_back(key.vectorized ? "VEC" : "SCALAR");
uint32_t wg_size = WEBGPU_MUL_MAT_VEC_WG_SIZE;
uint32_t outputs_per_wg = WEBGPU_MUL_MAT_VEC_FLOAT_OUTPUTS_PER_WG;
if (key.src0_type == GGML_TYPE_Q1_0) {
outputs_per_wg = WEBGPU_MUL_MAT_VEC_LEGACY_Q_OUTPUTS_PER_WG;
} else if (key.src0_type >= GGML_TYPE_Q2_K) {
outputs_per_wg = WEBGPU_MUL_MAT_VEC_K_Q_OUTPUTS_PER_WG;
} else if (key.src0_type >= GGML_TYPE_Q4_0) {
outputs_per_wg = WEBGPU_MUL_MAT_VEC_LEGACY_Q_OUTPUTS_PER_WG;
}
// variant suffix for src1 type
variant += std::string("_") + (context.src1->type == GGML_TYPE_F32 ? "f32" : "f16");
defines.push_back(std::string("WG_SIZE=") + std::to_string(wg_size));
defines.push_back(std::string("OUTPUTS_PER_WG=") + std::to_string(outputs_per_wg));
defines.push_back(context.supports_subgroups ? "USE_SUBGROUP_REDUCTION" : "USE_WORKGROUP_REDUCTION");
variant += context.supports_subgroups ? "_sg_reduce" : "_wg_reduce";
if (key.vectorized) {
variant += "_vectorized";
}
defines.push_back(std::string("N_EXPERTS=") + std::to_string(key.n_experts));
auto processed = preprocessor.preprocess(shader_src, defines);
auto decisions = std::make_shared<ggml_webgpu_mul_mat_vec_shader_decisions>();
decisions->wg_size = wg_size;
decisions->outputs_per_wg = outputs_per_wg;
webgpu_pipeline pipeline = ggml_webgpu_create_pipeline(device, processed, variant);
pipeline.context = decisions;
mul_mat_id_vec_pipelines[key] = pipeline;
return mul_mat_id_vec_pipelines[key];
}
webgpu_pipeline get_unary_pipeline(const ggml_webgpu_shader_lib_context & context) {
const bool is_unary = context.dst->op == GGML_OP_UNARY;
const int op = is_unary ? (int) ggml_get_unary_op(context.dst) : context.dst->op;
@@ -2800,6 +2978,72 @@ class ggml_webgpu_shader_lib {
return im2col_pipelines[key];
}
webgpu_pipeline get_upscale_pipeline(const ggml_webgpu_shader_lib_context & context) {
const uint32_t mode_flags = (uint32_t) ggml_get_op_params_i32(context.dst, 0);
const uint32_t base_mode = mode_flags & 0xFFu;
const bool antialias = (mode_flags & GGML_SCALE_FLAG_ANTIALIAS) != 0u;
ggml_webgpu_upscale_pipeline_key key = {};
key.input_type = context.src0->type;
key.output_type = context.dst->type;
key.base_mode = base_mode;
key.antialias = antialias;
auto it = upscale_pipelines.find(key);
if (it != upscale_pipelines.end()) {
return it->second;
}
std::vector<std::string> defines;
std::string variant = "upscale";
if (key.input_type == GGML_TYPE_F16) {
defines.push_back("SRC_F16");
variant += "_src_f16";
} else {
variant += "_src_f32";
}
if (key.output_type == GGML_TYPE_F16) {
defines.push_back("DST_F16");
variant += "_dst_f16";
} else {
variant += "_dst_f32";
}
switch (base_mode) {
case GGML_SCALE_MODE_NEAREST:
defines.push_back("NEAREST");
variant += "_nearest";
break;
case GGML_SCALE_MODE_BILINEAR:
defines.push_back("BILINEAR");
variant += "_bilinear";
break;
case GGML_SCALE_MODE_BICUBIC:
defines.push_back("BICUBIC");
variant += "_bicubic";
break;
default:
GGML_ABORT("Unsupported upscale mode");
}
if (antialias) {
defines.push_back("ANTIALIAS");
variant += "_aa";
}
defines.push_back(std::string("WG_SIZE=") + std::to_string(context.max_wg_size));
auto processed = preprocessor.preprocess(wgsl_upscale, defines);
auto decisions = std::make_shared<ggml_webgpu_generic_shader_decisions>();
decisions->wg_size = context.max_wg_size;
webgpu_pipeline pipeline = ggml_webgpu_create_pipeline(device, processed, variant);
pipeline.context = decisions;
upscale_pipelines[key] = pipeline;
return upscale_pipelines[key];
}
private:
static webgpu_pipeline ggml_webgpu_create_pipeline(wgpu::Device & device,
std::string shader_code,
+122 -2
View File
@@ -1404,7 +1404,6 @@ static webgpu_encoded_op ggml_webgpu_mul_mat(webgpu_context & ctx,
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1:
case GGML_TYPE_Q8_0:
case GGML_TYPE_Q8_1:
case GGML_TYPE_Q6_K:
case GGML_TYPE_Q4_K:
case GGML_TYPE_Q5_K:
@@ -1527,11 +1526,74 @@ static webgpu_encoded_op ggml_webgpu_mul_mat(webgpu_context & ctx,
return ggml_backend_webgpu_build(ctx, pipeline, params, entries, wg_x, wg_y);
}
static webgpu_encoded_op ggml_webgpu_mul_mat_id_vec(webgpu_context & ctx,
ggml_tensor * src0,
ggml_tensor * src1,
ggml_tensor * src2,
ggml_tensor * dst) {
const uint32_t param_n_expert = (uint32_t) src0->ne[2];
const uint32_t param_n_expert_used = (uint32_t) dst->ne[1];
ggml_webgpu_shader_lib_context shader_lib_ctx = {};
shader_lib_ctx.src0 = src0;
shader_lib_ctx.src1 = src1;
shader_lib_ctx.src2 = src2;
shader_lib_ctx.dst = dst;
shader_lib_ctx.supports_subgroups = ctx->global_ctx->capabilities.supports_subgroups;
shader_lib_ctx.max_wg_size = ctx->global_ctx->capabilities.limits.maxComputeInvocationsPerWorkgroup;
webgpu_pipeline pipeline = ctx->shader_lib->get_mul_mat_id_vec_pipeline(shader_lib_ctx);
std::vector<uint32_t> params = {
(uint32_t) (ggml_webgpu_tensor_misalignment(ctx, src0) / ggml_type_size(src0->type)),
(uint32_t) (ggml_webgpu_tensor_misalignment(ctx, src1) / ggml_type_size(src1->type)),
(uint32_t) (ggml_webgpu_tensor_misalignment(ctx, src2) / ggml_type_size(src2->type)),
(uint32_t) (ggml_webgpu_tensor_misalignment(ctx, dst) / ggml_type_size(dst->type)),
(uint32_t) src0->ne[0],
(uint32_t) src0->ne[1],
param_n_expert,
param_n_expert_used,
(uint32_t) src1->ne[1],
(uint32_t) (src0->nb[1] / ggml_type_size(src0->type)),
(uint32_t) (src1->nb[1] / ggml_type_size(src1->type)),
(uint32_t) (src0->nb[2] / ggml_type_size(src0->type)),
(uint32_t) (src1->nb[2] / ggml_type_size(src1->type)),
};
std::vector<wgpu::BindGroupEntry> entries = {
ggml_webgpu_make_bind_group_entry(0, ggml_webgpu_tensor_buf(src0), ggml_webgpu_tensor_align_offset(ctx, src0),
ggml_webgpu_tensor_binding_size(ctx, src0)),
ggml_webgpu_make_bind_group_entry(1, ggml_webgpu_tensor_buf(src1), ggml_webgpu_tensor_align_offset(ctx, src1),
ggml_webgpu_tensor_binding_size(ctx, src1)),
ggml_webgpu_make_bind_group_entry(2, ggml_webgpu_tensor_buf(src2), ggml_webgpu_tensor_align_offset(ctx, src2),
ggml_webgpu_tensor_binding_size(ctx, src2)),
ggml_webgpu_make_bind_group_entry(3, ggml_webgpu_tensor_buf(dst), ggml_webgpu_tensor_align_offset(ctx, dst),
ggml_webgpu_tensor_binding_size(ctx, dst)),
};
uint32_t wg_x = 1;
uint32_t wg_y = 1;
auto * decisions = static_cast<ggml_webgpu_mul_mat_vec_shader_decisions *>(pipeline.context.get());
const uint32_t max_wg_per_dim = ctx->global_ctx->capabilities.limits.maxComputeWorkgroupsPerDimension;
uint32_t output_groups = CEIL_DIV(dst->ne[0], decisions->outputs_per_wg);
uint32_t total_wg = output_groups * param_n_expert_used;
compute_2d_workgroups(total_wg, max_wg_per_dim, wg_x, wg_y);
return ggml_backend_webgpu_build(ctx, pipeline, params, entries, wg_x, wg_y);
}
static webgpu_encoded_op ggml_webgpu_mul_mat_id(webgpu_context & ctx,
ggml_tensor * src0,
ggml_tensor * src1,
ggml_tensor * src2,
ggml_tensor * dst) {
// we can use mat-vec fast path
if (dst->ne[2] == 1) {
return ggml_webgpu_mul_mat_id_vec(ctx, src0, src1, src2, dst);
}
ggml_webgpu_shader_lib_context shader_lib_ctx = {};
shader_lib_ctx.src0 = src0;
shader_lib_ctx.src1 = src1;
@@ -2762,6 +2824,49 @@ static bool ggml_webgpu_can_fuse_rms_norm_mul(const struct ggml_cgraph * cgraph,
return true;
}
static webgpu_encoded_op ggml_webgpu_upscale(webgpu_context ctx, ggml_tensor * src, ggml_tensor * dst) {
const uint32_t mode_flags = (uint32_t) ggml_get_op_params_i32(dst, 0);
std::vector<uint32_t> params = { (uint32_t) (ggml_webgpu_tensor_misalignment(ctx, src) / ggml_type_size(src->type)),
(uint32_t) (ggml_webgpu_tensor_misalignment(ctx, dst) / ggml_type_size(dst->type)),
(uint32_t) (src->nb[0] / ggml_type_size(src->type)),
(uint32_t) (src->nb[1] / ggml_type_size(src->type)),
(uint32_t) (src->nb[2] / ggml_type_size(src->type)),
(uint32_t) (src->nb[3] / ggml_type_size(src->type)),
(uint32_t) (dst->nb[0] / ggml_type_size(dst->type)),
(uint32_t) (dst->nb[1] / ggml_type_size(dst->type)),
(uint32_t) (dst->nb[2] / ggml_type_size(dst->type)),
(uint32_t) (dst->nb[3] / ggml_type_size(dst->type)),
(uint32_t) src->ne[0],
(uint32_t) src->ne[1],
(uint32_t) src->ne[2],
(uint32_t) src->ne[3],
(uint32_t) dst->ne[0],
(uint32_t) dst->ne[1],
(uint32_t) dst->ne[2],
(uint32_t) dst->ne[3],
mode_flags };
std::vector<wgpu::BindGroupEntry> entries = { ggml_webgpu_make_tensor_bind_group_entry(ctx, 0, src),
ggml_webgpu_make_tensor_bind_group_entry(ctx, 1, dst) };
ggml_webgpu_shader_lib_context shader_lib_ctx = {};
shader_lib_ctx.src0 = src;
shader_lib_ctx.dst = dst;
shader_lib_ctx.max_wg_size = ctx->global_ctx->capabilities.limits.maxComputeInvocationsPerWorkgroup;
webgpu_pipeline pipeline = ctx->shader_lib->get_upscale_pipeline(shader_lib_ctx);
auto * decisions = static_cast<ggml_webgpu_generic_shader_decisions *>(pipeline.context.get());
uint32_t total_wg = CEIL_DIV((uint32_t) ggml_nelements(dst), decisions->wg_size);
uint32_t wg_x = std::min(ctx->global_ctx->capabilities.limits.maxComputeWorkgroupsPerDimension, total_wg);
uint32_t wg_y = CEIL_DIV(total_wg, wg_x);
return ggml_backend_webgpu_build(ctx, pipeline, params, entries, wg_x, wg_y);
}
// Returns the encoded command, or std::nullopt if the operation is a no-op
static std::optional<webgpu_encoded_op> ggml_webgpu_encode(webgpu_context ctx,
ggml_cgraph * cgraph,
@@ -2869,6 +2974,8 @@ static std::optional<webgpu_encoded_op> ggml_webgpu_encode(webgpu_context ctx,
return ggml_webgpu_conv_2d(ctx, src0, src1, node);
case GGML_OP_IM2COL:
return ggml_webgpu_im2col(ctx, src0, src1, node);
case GGML_OP_UPSCALE:
return ggml_webgpu_upscale(ctx, src0, node);
default:
return std::nullopt;
}
@@ -3107,8 +3214,8 @@ static ggml_backend_i ggml_backend_webgpu_i = {
/* .free = */ ggml_backend_webgpu_free,
/* .set_tensor_async = */ ggml_backend_webgpu_set_tensor_async,
/* .get_tensor_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .set_tensor_2d_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .cpy_tensor_async = */ NULL,
/* .synchronize = */ ggml_backend_webgpu_synchronize,
/* .graph_plan_create = */ NULL,
@@ -3879,6 +3986,15 @@ static bool ggml_backend_webgpu_device_supports_op(ggml_backend_dev_t dev, const
case GGML_TYPE_Q4_K:
case GGML_TYPE_Q5_K:
case GGML_TYPE_Q6_K:
case GGML_TYPE_IQ1_S:
case GGML_TYPE_IQ1_M:
case GGML_TYPE_IQ2_XXS:
case GGML_TYPE_IQ2_XS:
case GGML_TYPE_IQ2_S:
case GGML_TYPE_IQ3_XXS:
case GGML_TYPE_IQ3_S:
case GGML_TYPE_IQ4_NL:
case GGML_TYPE_IQ4_XS:
supports_op = true;
break;
default:
@@ -4092,6 +4208,10 @@ static bool ggml_backend_webgpu_device_supports_op(ggml_backend_dev_t dev, const
case GGML_OP_SUM_ROWS:
supports_op = op->type == GGML_TYPE_F32 && src0->type == op->type && ggml_is_contiguous_rows(src0);
break;
case GGML_OP_UPSCALE:
supports_op = (op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16) &&
(src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16);
break;
default:
break;
}
@@ -0,0 +1,154 @@
#ifdef USE_SUBGROUP_REDUCTION
enable subgroups;
#endif
enable f16;
#define DECLARE_BYTE_LOADERS_SRC0
#include "common_decls.tmpl"
#include "mul_mat_vec_acc.tmpl"
struct MulMatIdVecParams {
offset_src0: u32,
offset_src1: u32,
offset_ids: u32,
offset_dst: u32,
k: u32,
m: u32,
n_expert: u32,
n_expert_used: u32,
b_ne1: u32,
stride_01: u32,
stride_11: u32,
stride_02: u32,
stride_12: u32,
};
@group(0) @binding(0) var<storage, read_write> src0: array<SRC0_TYPE>; // [cols, rows, n_expert]
@group(0) @binding(1) var<storage, read_write> src1: array<SRC1_TYPE>; // [cols, b_ne1, n_tokens(1)]
@group(0) @binding(2) var<storage, read_write> ids: array<u32>; // [n_experd_used, n_tokens(1)]
@group(0) @binding(3) var<storage, read_write> dst: array<f32>; // [rows, n_expert_used, n_tokens(1)]
// "mul_mat_vec_acc.tmpl" requires params.k, params.m, params.stride_01
@group(0) @binding(4) var<uniform> params: MulMatIdVecParams;
// Flattened as [row][thread] to keep each row's reduction contiguous in memory.
var<workgroup> partial_sums: array<f32, OUTPUTS_PER_WG * WG_SIZE>;
fn partial_index(row: u32, thread: u32) -> u32 {
return row * WG_SIZE + thread;
}
var<workgroup> gathered_count_ids: array<u32, N_EXPERTS>;
var<workgroup> gathered_expert_used: array<u32, N_EXPERTS>;
@compute @workgroup_size(WG_SIZE)
fn main(
@builtin(local_invocation_id) local_id: vec3<u32>,
@builtin(workgroup_id) wg_id: vec3<u32>,
@builtin(num_workgroups) num_wg: vec3<u32>
#ifdef USE_SUBGROUP_REDUCTION
, @builtin(subgroup_id) subgroup_id: u32,
@builtin(subgroup_invocation_id) subgroup_invocation_id: u32,
@builtin(num_subgroups) num_subgroups: u32,
@builtin(subgroup_size) subgroup_size: u32
#endif
) {
let thread_id = local_id.x;
for (var i = thread_id;i < params.n_expert;i += WG_SIZE) {
gathered_count_ids[i] = 0;
}
workgroupBarrier();
// gather the selected experts for the target token.
for (var col = thread_id;col < params.n_expert_used;col += WG_SIZE) {
let expert = ids[params.offset_ids + col];
gathered_count_ids[expert] = 1;
gathered_expert_used[expert] = col;
}
workgroupBarrier();
let output_groups:u32 = (params.m + OUTPUTS_PER_WG - 1u) / OUTPUTS_PER_WG;
let wg_linear = wg_id.y * num_wg.x + wg_id.x;
var own_expert:u32 = 0;
var wg_in_batch:u32 = 0;
var wg_sum:u32 = 0;
for (var i = 0u;i < params.n_expert;i += 1) {
let wg_vec_count = gathered_count_ids[i]; // 1 or 0
let wg_per_matrix = output_groups * wg_vec_count;
if (wg_sum <= wg_linear && wg_linear < wg_sum + wg_per_matrix) {
own_expert = i;
wg_in_batch = wg_linear - wg_sum;
break;
}
wg_sum += wg_per_matrix;
}
let row_base = (wg_linear % output_groups) * OUTPUTS_PER_WG;
let dst1_stride = params.m;
let src0_batch_offset = params.offset_src0 + own_expert * params.stride_02;
let src1_idx_base = params.offset_src1 + (gathered_expert_used[own_expert] % params.b_ne1) * params.stride_11;
let dst_idx_base = params.offset_dst + gathered_expert_used[own_expert] * dst1_stride + row_base;
let acc = accumulate_vec_dot(thread_id, row_base, src0_batch_offset, src1_idx_base);
#ifdef USE_SUBGROUP_REDUCTION
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
let subgroup_total = subgroupAdd(acc[row]);
if (subgroup_invocation_id == 0u) {
partial_sums[partial_index(row, subgroup_id)] = subgroup_total;
}
}
workgroupBarrier();
for (var row = subgroup_id; (row < OUTPUTS_PER_WG) && (row_base + row < params.m); row += num_subgroups) {
let output_row = row_base + row;
var row_acc = 0.0f;
for (var k = subgroup_invocation_id; k < num_subgroups; k += subgroup_size) {
row_acc += partial_sums[partial_index(row, k)];
}
let row_total = subgroupAdd(row_acc);
if (subgroup_invocation_id == 0) {
dst[dst_idx_base + row] = row_total;
}
}
#endif
#ifdef USE_WORKGROUP_REDUCTION
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
partial_sums[partial_index(row, thread_id)] = acc[row];
}
workgroupBarrier();
var stride:u32 = WG_SIZE / 2u;
while (stride > 0) {
if (thread_id < stride) {
for (var row = 0u; row < OUTPUTS_PER_WG; row++) {
partial_sums[partial_index(row, thread_id)] += partial_sums[partial_index(row, thread_id + stride)];
}
}
workgroupBarrier();
stride = stride / 2;
}
if (thread_id < OUTPUTS_PER_WG) {
let output_row = row_base + thread_id;
if (output_row < params.m) {
dst[dst_idx_base + thread_id] = partial_sums[partial_index(thread_id, 0)];
}
}
#endif
}
File diff suppressed because it is too large Load Diff
File diff suppressed because it is too large Load Diff
@@ -0,0 +1,240 @@
#if defined(SRC_F16) || defined(DST_F16)
enable f16;
#endif
#ifdef SRC_F16
#define SRC_TYPE f16
#else
#define SRC_TYPE f32
#endif
#ifdef DST_F16
#define DST_TYPE f16
#else
#define DST_TYPE f32
#endif
@group(0) @binding(0)
var<storage, read_write> input: array<SRC_TYPE>;
@group(0) @binding(1)
var<storage, read_write> output: array<DST_TYPE>;
struct Params {
offset_i: u32,
offset_o: u32,
// element strides
si0: u32, si1: u32, si2: u32, si3: u32,
so0: u32, so1: u32, so2: u32, so3: u32,
src_w: u32,
src_h: u32,
src_z: u32,
src_n: u32,
dst_w: u32,
dst_h: u32,
dst_z: u32,
dst_n: u32,
mode_flags: u32,
};
@group(0) @binding(2)
var<uniform> params: Params;
const GGML_SCALE_FLAG_ALIGN_CORNERS: u32 = 1u << 8u;
fn get_clamped_input(x: i32, y: i32, z: u32, n: u32) -> f32 {
let cx = u32(clamp(x, 0, i32(params.src_w) - 1));
let cy = u32(clamp(y, 0, i32(params.src_h) - 1));
let i = params.offset_i + cx * params.si0 + cy * params.si1 + z * params.si2 + n * params.si3;
return f32(input[i]);
}
fn cubic_weight(t: f32, a: f32) -> f32 {
let at = abs(t);
if (at <= 1.0) {
return (a + 2.0) * at * at * at - (a + 3.0) * at * at + 1.0;
} else if (at <= 2.0) {
return a * at * at * at - 5.0 * a * at * at + 8.0 * a * at - 4.0 * a;
} else {
return 0.0;
}
}
@compute @workgroup_size(WG_SIZE)
fn main(
@builtin(global_invocation_id) gid: vec3<u32>,
@builtin(num_workgroups) num_wg: vec3<u32>
) {
let i_out = gid.x + (num_wg.x * u32(WG_SIZE)) * gid.y;
let total = params.dst_w * params.dst_h * params.dst_z * params.dst_n;
if (i_out >= total) {
return;
}
// decode (x, y, z, n)
var i = i_out;
let x_dst = i % params.dst_w;
i = i / params.dst_w;
let y_dst = i % params.dst_h;
i = i / params.dst_h;
let z_dst = i % params.dst_z;
let n_dst = i / params.dst_z;
// scale factors
var sf0 = f32(params.dst_w) / f32(params.src_w);
var sf1 = f32(params.dst_h) / f32(params.src_h);
var sf2 = f32(params.dst_z) / f32(params.src_z);
var sf3 = f32(params.dst_n) / f32(params.src_n);
let align_corners = (params.mode_flags & GGML_SCALE_FLAG_ALIGN_CORNERS) != 0;
// pixel_offset: 0.5 for half-pixel-center (default), 0.0 for align_corners
var pixel_offset = 0.5;
if (align_corners) {
pixel_offset = 0.0;
if (params.dst_w > 1 && params.src_w > 1) {
sf0 = f32(params.dst_w - 1) / f32(params.src_w - 1);
}
if (params.dst_h > 1 && params.src_h > 1) {
sf1 = f32(params.dst_h - 1) / f32(params.src_h - 1);
}
}
let z_src = min(params.src_z - 1, u32(floor(f32(z_dst) / sf2)));
let n_src = min(params.src_n - 1, u32(floor(f32(n_dst) / sf3)));
var result = 0.0;
#if defined(NEAREST)
let x_src = min(params.src_w - 1, u32(floor(f32(x_dst) / sf0)));
let y_src = min(params.src_h - 1, u32(floor(f32(y_dst) / sf1)));
result = get_clamped_input(i32(x_src), i32(y_src), z_src, n_src);
#elif defined(BILINEAR)
#if defined(ANTIALIAS)
// Antialiased bilinear: triangle filter over a variable support region.
let support0 = max(1.0f / sf0, 1.0f);
let support1 = max(1.0f / sf1, 1.0f);
let invscale0 = 1.0 / support0;
let invscale1 = 1.0 / support1;
let fx = (f32(x_dst) + pixel_offset) / sf0;
let fy = (f32(y_dst) + pixel_offset) / sf1;
let x_min = max(i32(fx - support0 + pixel_offset), 0);
let y_min = max(i32(fy - support1 + pixel_offset), 0);
let x_max = min(i32(fx + support0 + pixel_offset), i32(params.src_w));
let y_max = min(i32(fy + support1 + pixel_offset), i32(params.src_h));
var weighted_sum = 0.0;
var total_weight = 0.0;
for (var x = x_min; x < x_max; x += 1) {
let wx = max(1.0 - abs(f32(x) - fx + pixel_offset) * invscale0, 0.0);
for (var y = y_min; y < y_max; y += 1) {
let wy = max(1.0 - abs(f32(y) - fy + pixel_offset) * invscale1, 0.0);
let w = wx * wy;
if (w > 0.0) {
weighted_sum += get_clamped_input(x, y, z_src, n_src) * w;
total_weight += w;
}
}
}
if (total_weight > 0.0) {
result = weighted_sum / total_weight;
}
#else
let fx = (f32(x_dst) + pixel_offset) / sf0 - pixel_offset;
let fy = (f32(y_dst) + pixel_offset) / sf1 - pixel_offset;
let x0 = i32(floor(fx));
let y0 = i32(floor(fy));
let dx = clamp(fx - f32(x0), 0.0, 1.0);
let dy = clamp(fy - f32(y0), 0.0, 1.0);
let a = get_clamped_input(x0, y0, z_src, n_src);
let b = get_clamped_input(x0 + 1, y0, z_src, n_src);
let c = get_clamped_input(x0, y0 + 1, z_src, n_src);
let d = get_clamped_input(x0 + 1, y0 + 1, z_src, n_src);
let wa = (1.0 - dx) * (1.0 - dy);
let wb = dx * (1.0 - dy);
let wc = (1.0 - dx) * dy;
let wd = dx * dy;
result = a * wa + b * wb + c * wc + d * wd;
#endif
#elif defined(BICUBIC)
// bicubic convolution with alpha = -0.75 (PyTorch default)
let alpha = -0.75;
let fx = (f32(x_dst) + pixel_offset) / sf0 - pixel_offset;
let fy = (f32(y_dst) + pixel_offset) / sf1 - pixel_offset;
let x0 = i32(floor(fx));
let y0 = i32(floor(fy));
let dx = fx - f32(x0);
let dy = fy - f32(y0);
// horizontal weights for offsets -1, 0, 1, 2
let wx0 = cubic_weight(dx + 1.0, alpha);
let wx1 = cubic_weight(dx, alpha);
let wx2 = cubic_weight(1.0 - dx, alpha);
let wx3 = cubic_weight(2.0 - dx, alpha);
// vertical weights for offsets -1, 0, 1, 2
let wy0 = cubic_weight(dy + 1.0, alpha);
let wy1 = cubic_weight(dy, alpha);
let wy2 = cubic_weight(1.0 - dy, alpha);
let wy3 = cubic_weight(2.0 - dy, alpha);
// intermediate horizontal interpolation for 4x4 grid of pixels
// x0-1, x0, x0+1, x0+2, y0-1
let p0 = get_clamped_input(x0 - 1, y0 - 1, z_src, n_src);
let p1 = get_clamped_input(x0, y0 - 1, z_src, n_src);
let p2 = get_clamped_input(x0 + 1, y0 - 1, z_src, n_src);
let p3 = get_clamped_input(x0 + 2, y0 - 1, z_src, n_src);
let row0 = p0 * wx0 + p1 * wx1 + p2 * wx2 + p3 * wx3;
// x0-1, x0, x0+1, x0+2, y0
let q0 = get_clamped_input(x0 - 1, y0, z_src, n_src);
let q1 = get_clamped_input(x0, y0, z_src, n_src);
let q2 = get_clamped_input(x0 + 1, y0, z_src, n_src);
let q3 = get_clamped_input(x0 + 2, y0, z_src, n_src);
let row1 = q0 * wx0 + q1 * wx1 + q2 * wx2 + q3 * wx3;
// x0-1, x0, x0+1, x0+2, y0+1
let r0 = get_clamped_input(x0 - 1, y0 + 1, z_src, n_src);
let r1 = get_clamped_input(x0, y0 + 1, z_src, n_src);
let r2 = get_clamped_input(x0 + 1, y0 + 1, z_src, n_src);
let r3 = get_clamped_input(x0 + 2, y0 + 1, z_src, n_src);
let row2 = r0 * wx0 + r1 * wx1 + r2 * wx2 + r3 * wx3;
// x0-1, x0, x0+1, x0+2, y0+2
let s0 = get_clamped_input(x0 - 1, y0 + 2, z_src, n_src);
let s1 = get_clamped_input(x0, y0 + 2, z_src, n_src);
let s2 = get_clamped_input(x0 + 1, y0 + 2, z_src, n_src);
let s3 = get_clamped_input(x0 + 2, y0 + 2, z_src, n_src);
let row3 = s0 * wx0 + s1 * wx1 + s2 * wx2 + s3 * wx3;
// final vertical interpolation
result = row0 * wy0 + row1 * wy1 + row2 * wy2 + row3 * wy3;
#endif
let dst_idx = params.offset_o + x_dst * params.so0 + y_dst * params.so1 + z_dst * params.so2 + n_dst * params.so3;
output[dst_idx] = DST_TYPE(result);
}
+1 -1
View File
@@ -423,8 +423,8 @@ static ggml_backend_i ggml_backend_zdnn_i = {
/* .free = */ ggml_backend_zdnn_free,
/* .set_tensor_async = */ NULL,
/* .get_tensor_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .set_tensor_2d_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .cpy_tensor_async = */ NULL,
/* .synchronize = */ NULL,
/* .graph_plan_create = */ NULL,
+1 -1
View File
@@ -407,8 +407,8 @@ static struct ggml_backend_i ggml_backend_zendnn_i = {
/* .free = */ ggml_backend_zendnn_free,
/* .set_tensor_async = */ NULL,
/* .get_tensor_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .set_tensor_2d_async = */ NULL,
/* .get_tensor_2d_async = */ NULL,
/* .cpy_tensor_async = */ NULL,
/* .synchronize = */ NULL,
/* .graph_plan_create = */ NULL,
+2 -1
View File
@@ -20,6 +20,7 @@ from PySide6.QtCore import Qt, QRect, QSize
from jinja2 import TemplateSyntaxError
from jinja2.sandbox import ImmutableSandboxedEnvironment
from datetime import datetime
from typing import Callable
def format_template_content(template_content):
@@ -395,7 +396,7 @@ class JinjaTester(QMainWindow):
ensure_ascii=ensure_ascii,
)
)
env.globals["strftime_now"] = lambda format: datetime.now().strftime(format) # ty: ignore[invalid-assignment]
env.globals["strftime_now"]: Callable[[str], str] = lambda format: datetime.now().strftime(format)
env.globals["raise_exception"] = raise_exception # ty: ignore[invalid-assignment]
try:
template = env.from_string(template_str)
+1 -1
View File
@@ -5,7 +5,7 @@ import os
import sys
import subprocess
HTTPLIB_VERSION = "refs/tags/v0.43.1"
HTTPLIB_VERSION = "refs/tags/v0.43.2"
vendor = {
"https://github.com/nlohmann/json/releases/latest/download/json.hpp": "vendor/nlohmann/json.hpp",
+58
View File
@@ -0,0 +1,58 @@
#!/usr/bin/env bash
# initialize a new worktree from a branch name:
#
# - creates a new branch from current HEAD
# - creates a new worktree in a parent folder, suffixed with the branch name
#
# sample usage:
# ./scripts/wc2wt.sh gg/new-feature-foo-bar
# ./scripts/wc2wt.sh gg/new-feature-foo-bar opencode
# ./scripts/wc2wt.sh gg/new-feature-foo-bar "cmake -B build && cmake --build build"
# ./scripts/wc2wt.sh gg/new-feature-foo-bar "bash -l"
function usage() {
echo "usage: $0 <branch_name> [cmd]"
exit 1
}
# check we are in the right directory
if [[ ! -f "scripts/wc2wt.sh" ]]; then
echo "error: this script must be run from the root of the repository"
exit 1
fi
if [[ $# -lt 1 || $# -gt 2 ]]; then
usage
fi
BRANCH=$1
if [[ -z "$BRANCH" ]]; then
echo "error: branch name must not be empty"
exit 1
fi
dir=$(basename $(pwd))
# sanitize branch name for directory name (replace / with -)
dir_suffix=$(echo "$BRANCH" | tr '/' '-')
git worktree add -b "$BRANCH" "../$dir-$dir_suffix" HEAD
og_path=$(pwd)
wt_path=$(cd "../$dir-$dir_suffix" && pwd)
echo "git worktree created in $wt_path"
cd "$wt_path"
# pi agent setup in the worktree
if [[ -f "$og_path/.pi/SYSTEM.md" && ! -f ".pi/SYSTEM.md" ]]; then
mkdir -p .pi
ln -sfn "$og_path/.pi/SYSTEM.md" .pi/SYSTEM.md
fi
if [[ $# -eq 2 ]]; then
echo "executing: $2"
eval "$2"
fi
+10 -2
View File
@@ -40,6 +40,14 @@
#include <TargetConditionals.h>
#endif
#ifdef _WIN32
# define llama_mmap_ftell _ftelli64
# define llama_mmap_fseek _fseeki64
#else
# define llama_mmap_ftell ftello
# define llama_mmap_fseek fseeko
#endif
// TODO: consider moving to llama-impl.h if needed in more places
#if defined(_WIN32)
static std::string llama_format_win_err(DWORD err) {
@@ -226,7 +234,7 @@ struct llama_file::impl {
size_t tell() const {
if (fd == -1) {
long ret = std::ftell(fp);
off_t ret = llama_mmap_ftell(fp);
if (ret == -1) {
throw std::runtime_error(format("ftell error: %s", strerror(errno)));
}
@@ -244,7 +252,7 @@ struct llama_file::impl {
void seek(size_t offset, int whence) const {
off_t ret = 0;
if (fd == -1) {
ret = std::fseek(fp, (long) offset, whence);
ret = llama_mmap_fseek(fp, offset, whence);
} else {
ret = lseek(fd, offset, whence);
}
+2 -2
View File
@@ -683,9 +683,9 @@ static ggml_type llama_tensor_get_type(quantize_state_impl & qs, const llama_mod
LLAMA_LOG_WARN("%s: %-36s - applying manual override: %s -> %s\n",
__func__, tensor_name.c_str(), ggml_type_name(new_type), ggml_type_name(qtype));
new_type = qtype;
manual = true;
break;
}
manual = true;
break;
}
}
}
+31 -15
View File
@@ -6244,11 +6244,12 @@ struct test_flash_attn_ext : public test_case {
const float logit_softcap; // Gemma 2
const ggml_prec prec;
const ggml_type type_KV;
const ggml_type type_K;
const ggml_type type_V;
std::array<int32_t, 4> permute;
std::string vars() override {
return VARS_TO_STR13(hsk, hsv, nh, nr23, kv, nb, mask, sinks, max_bias, logit_softcap, prec, type_KV, permute);
return VARS_TO_STR14(hsk, hsv, nh, nr23, kv, nb, mask, sinks, max_bias, logit_softcap, prec, type_K, type_V, permute);
}
double max_nmse_err() override {
@@ -6264,12 +6265,13 @@ struct test_flash_attn_ext : public test_case {
test_flash_attn_ext(int64_t hsk = 128, int64_t hsv = 128, int64_t nh = 32, std::array<int64_t, 2> nr23 = {1, 1}, int64_t kv = 96, int64_t nb = 8,
bool mask = true, bool sinks = false, float max_bias = 0.0f, float logit_softcap = 0.0f, ggml_prec prec = GGML_PREC_F32,
ggml_type type_KV = GGML_TYPE_F16, std::array<int32_t, 4> permute = {0, 1, 2, 3})
: hsk(hsk), hsv(hsv), nh(nh), nr23(nr23), kv(kv), nb(nb), mask(mask), sinks(sinks), max_bias(max_bias), logit_softcap(logit_softcap), prec(prec), type_KV(type_KV), permute(permute) {}
ggml_type type_K = GGML_TYPE_F16, ggml_type type_V = GGML_TYPE_F16, std::array<int32_t, 4> permute = {0, 1, 2, 3})
: hsk(hsk), hsv(hsv), nh(nh), nr23(nr23), kv(kv), nb(nb), mask(mask), sinks(sinks), max_bias(max_bias), logit_softcap(logit_softcap), prec(prec),
type_K(type_K), type_V(type_V), permute(permute) {}
ggml_tensor * build_graph(ggml_context * ctx) override {
const int64_t hsk_padded = GGML_PAD(hsk, ggml_blck_size(type_KV));
const int64_t hsv_padded = GGML_PAD(hsv, ggml_blck_size(type_KV));
const int64_t hsk_padded = GGML_PAD(hsk, ggml_blck_size(type_K));
const int64_t hsv_padded = GGML_PAD(hsv, ggml_blck_size(type_V));
auto const &create_permuted = [&](ggml_type type, int64_t ne0, int64_t ne1, int64_t ne2, int64_t ne3, bool is_view) -> ggml_tensor * {
int64_t ne[4] = {ne0, ne1, ne2, ne3};
@@ -6293,11 +6295,11 @@ struct test_flash_attn_ext : public test_case {
ggml_tensor * q = create_permuted(GGML_TYPE_F32, hsk_padded, nb, nh*nr23[0], nr23[1], false);
ggml_set_name(q, "q");
ggml_tensor * k = create_permuted(type_KV, hsk_padded, kv, nh, nr23[1], true); // the K tensor is usually a view of the K cache
ggml_tensor * k = create_permuted(type_K, hsk_padded, kv, nh, nr23[1], true); // the K tensor is usually a view of the K cache
ggml_set_name(k, "k");
ggml_tensor * v = nullptr;
if (hsk_padded == 576 && hsv_padded == 512) {
if (type_K == type_V && hsk_padded == 576 && hsv_padded == 512) {
// TODO: this branch should become a separate test case parameter instead of hardcoding this for these head shapes
// in this branch, the V cache is sub-view of the K cache. this is used by some MLA-based models
@@ -6307,7 +6309,7 @@ struct test_flash_attn_ext : public test_case {
// - https://github.com/ggml-org/llama.cpp/pull/18986
v = ggml_view_4d(ctx, k, hsv_padded, kv, nh, nr23[1], k->nb[1], k->nb[2], k->nb[3], 0);
} else {
v = create_permuted(type_KV, hsv_padded, kv, nh, nr23[1], true); // the V tensor is usually a view of the V cache
v = create_permuted(type_V, hsv_padded, kv, nh, nr23[1], true); // the V tensor is usually a view of the V cache
}
ggml_set_name(v, "v");
@@ -8707,11 +8709,11 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
for (ggml_type type_KV : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_BF16, GGML_TYPE_Q8_0, GGML_TYPE_Q5_1, GGML_TYPE_Q5_0, GGML_TYPE_Q4_1, GGML_TYPE_Q4_0, GGML_TYPE_IQ4_NL}) {
if (type_KV != GGML_TYPE_F16 && hsk != 64 && hsk != 72) continue;
test_cases.emplace_back(new test_flash_attn_ext(
hsk, hsv, nh, {nr2, nr3}, kv, nb, mask, sinks, max_bias, logit_softcap, prec, type_KV));
hsk, hsv, nh, {nr2, nr3}, kv, nb, mask, sinks, max_bias, logit_softcap, prec, type_KV, type_KV));
// run fewer test cases permuted
if (mask == true && max_bias == 0.0f && logit_softcap == 0 && kv == 512) {
test_cases.emplace_back(new test_flash_attn_ext(
hsk, hsv, nh, {nr2, nr3}, kv, nb, mask, sinks, max_bias, logit_softcap, prec, type_KV, {0, 2, 1, 3}));
hsk, hsv, nh, {nr2, nr3}, kv, nb, mask, sinks, max_bias, logit_softcap, prec, type_KV, type_KV, {0, 2, 1, 3}));
}
}
}
@@ -8727,6 +8729,16 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
}
}
// mixed quant and Q1_0 test cases
test_cases.emplace_back(new test_flash_attn_ext(64, 64, 4, {1, 1}, 128, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q4_0));
test_cases.emplace_back(new test_flash_attn_ext(64, 64, 4, {1, 1}, 128, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q4_0, GGML_TYPE_F16));
test_cases.emplace_back(new test_flash_attn_ext(72, 72, 4, {1, 1}, 96, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0));
test_cases.emplace_back(new test_flash_attn_ext(64, 64, 4, {1, 1}, 96, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F32));
test_cases.emplace_back(new test_flash_attn_ext(128, 128, 4, {1, 1}, 96, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q1_0, GGML_TYPE_Q1_0));
test_cases.emplace_back(new test_flash_attn_ext(128, 64, 4, {1, 1}, 128, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q1_0, GGML_TYPE_Q4_0));
test_cases.emplace_back(new test_flash_attn_ext(64, 128, 4, {1, 1}, 128, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q4_0, GGML_TYPE_Q1_0));
test_cases.emplace_back(new test_flash_attn_ext(128, 64, 4, {1, 1}, 64, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q1_0, GGML_TYPE_F16));
test_cases.emplace_back(new test_cross_entropy_loss (GGML_TYPE_F32, { 10, 5, 4, 3}));
test_cases.emplace_back(new test_cross_entropy_loss (GGML_TYPE_F32, {30000, 1, 1, 1}));
test_cases.emplace_back(new test_cross_entropy_loss_back(GGML_TYPE_F32, { 10, 5, 4, 3}));
@@ -8971,15 +8983,19 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
}
// Qwen3-VL-8B https://github.com/ggml-org/llama.cpp/issues/17012
test_cases.emplace_back(new test_flash_attn_ext(72, 72, 16, {1, 1}, 5776, 5776, false, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16));
test_cases.emplace_back(new test_flash_attn_ext(72, 72, 16, {1, 1}, 5776, 5776, false, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16));
test_cases.emplace_back(new test_flash_attn_ext(64, 64, 8, {8, 1}, 7680, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16));
test_cases.emplace_back(new test_flash_attn_ext(64, 64, 8, {8, 1}, 7680, 4, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16));
test_cases.emplace_back(new test_flash_attn_ext(64, 64, 8, {8, 1}, 7680, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16));
test_cases.emplace_back(new test_flash_attn_ext(64, 64, 8, {8, 1}, 7680, 4, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16));
test_cases.emplace_back(new test_flash_attn_ext(64, 64, 8, {8, 1}, 7680, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q4_0, GGML_TYPE_Q4_0));
test_cases.emplace_back(new test_flash_attn_ext(64, 64, 8, {8, 1}, 7680, 512, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q4_0, GGML_TYPE_Q4_0));
test_cases.emplace_back(new test_flash_attn_ext(64, 64, 8, {8, 1}, 7680, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0));
test_cases.emplace_back(new test_flash_attn_ext(64, 64, 8, {8, 1}, 7680, 512, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0));
for (int kv : { 4096, 8192, 16384, }) {
for (int hs : { 64, 128, }) {
for (int nr : { 1, 4, }) {
test_cases.emplace_back(new test_flash_attn_ext(hs, hs, 8, {nr, 1}, kv, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16));
test_cases.emplace_back(new test_flash_attn_ext(hs, hs, 8, {nr, 1}, kv, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16));
}
}
}
File diff suppressed because one or more lines are too long
+4091 -4084
View File
File diff suppressed because it is too large Load Diff
@@ -1,5 +1,5 @@
<script lang="ts">
import * as Tooltip from '../src/lib/components/ui/tooltip';
import * as Tooltip from '../../src/lib/components/ui/tooltip';
interface Props {
children: any;
+2 -2
View File
@@ -1,7 +1,7 @@
import type { Preview } from '@storybook/sveltekit';
import '../src/app.css';
import ModeWatcherDecorator from './ModeWatcherDecorator.svelte';
import TooltipProviderDecorator from './TooltipProviderDecorator.svelte';
import ModeWatcherDecorator from './decorators/ModeWatcherDecorator.svelte';
import TooltipProviderDecorator from './decorators/TooltipProviderDecorator.svelte';
const preview: Preview = {
parameters: {
+3 -3
View File
@@ -3640,9 +3640,9 @@
}
},
"node_modules/bits-ui": {
"version": "2.17.3",
"resolved": "https://registry.npmjs.org/bits-ui/-/bits-ui-2.17.3.tgz",
"integrity": "sha512-Bef41uY9U2jaBJHPhcPvmBNkGec5Wx2z6eioDsTmsaR2vH4QoaOcPi75gzCG3+/2TNr6v/qBwzgWNPYCxNtrEA==",
"version": "2.18.0",
"resolved": "https://registry.npmjs.org/bits-ui/-/bits-ui-2.18.0.tgz",
"integrity": "sha512-GLOBZRVy3hxNHIQ2MpD/+5aK9KcBFZRhUJtZ1UDABXdlVR4K6zFpgt4T+Rwuhf2sQzlc6yK1q/DprHPjwT4Pjw==",
"dev": true,
"license": "MIT",
"dependencies": {
-2
View File
@@ -28,7 +28,6 @@ import type {
ApiRouterModelsUnloadResponse,
// Chat types
ChatAttachmentDisplayItem,
ChatAttachmentPreviewItem,
ChatMessageType,
ChatRole,
ChatUploadedFile,
@@ -92,7 +91,6 @@ declare global {
ApiRouterModelsUnloadResponse,
// Chat types
ChatAttachmentDisplayItem,
ChatAttachmentPreviewItem,
ChatMessagePromptProgress,
ChatMessageSiblingInfo,
ChatMessageTimings,
@@ -1,3 +1,5 @@
import { isElementInViewport } from '$lib/utils/viewport';
/**
* Svelte action that fades in an element when it enters the viewport.
* Uses IntersectionObserver for efficient viewport detection.
@@ -12,17 +14,8 @@ export function fadeInView(
) {
const { duration = 300, y = 0, skipIfVisible = false } = options;
if (skipIfVisible) {
const rect = node.getBoundingClientRect();
const isAlreadyVisible =
rect.top < window.innerHeight &&
rect.bottom > 0 &&
rect.left < window.innerWidth &&
rect.right > 0;
if (isAlreadyVisible) {
return;
}
if (skipIfVisible && isElementInViewport(node)) {
return;
}
node.style.opacity = '0';
@@ -0,0 +1,11 @@
---
name: app
description: Opinionated app components building on top of ./ui primitives
---
- Can include business logic and state management
- Can include data fetching and caching logic
- Should use original spelling for HTML-native events and `camelCase` for custom events
- Props and markup attributes should be listed alphabetically
- Use JS Objects and Arrays for CSS classes and styles when they are dynamic
- Whenever there can be repetition in the component's markup, if it's too small to be decoupled as a separate component — use Svelte 5's `{#snippet}` + `{@render}`
@@ -5,15 +5,16 @@
import { TooltipSide } from '$lib/enums';
interface Props {
icon: Component;
tooltip: string;
variant?: ButtonVariant;
size?: ButtonSize;
iconSize?: string;
ariaLabel?: string;
class?: string;
disabled?: boolean;
icon: Component;
iconSize?: string;
onclick: (e?: MouseEvent) => void;
'aria-label'?: string;
size?: ButtonSize;
stopPropagationOnClick?: boolean;
tooltip: string;
variant?: ButtonVariant;
tooltipSide?: TooltipSide;
}
@@ -26,8 +27,9 @@
disabled = false,
iconSize = 'h-3 w-3',
tooltipSide = TooltipSide.TOP,
stopPropagationOnClick = false,
onclick,
'aria-label': ariaLabel
ariaLabel
}: Props = $props();
</script>
@@ -37,13 +39,18 @@
{variant}
{size}
{disabled}
{onclick}
onclick={(e: MouseEvent) => {
if (stopPropagationOnClick) e.stopPropagation();
onclick?.(e);
}}
class="h-6 w-6 p-0 {className} flex hover:bg-transparent data-[state=open]:bg-transparent!"
aria-label={ariaLabel || tooltip}
>
{@const IconComponent = icon}
<IconComponent class={iconSize} />
{#if icon}
{@const IconComponent = icon}
<IconComponent class={iconSize} />
{/if}
</Button>
</Tooltip.Trigger>
@@ -1,18 +1,17 @@
<script lang="ts">
import { Copy } from '@lucide/svelte';
import { copyToClipboard } from '$lib/utils';
import ActionIcon from './ActionIcon.svelte';
interface Props {
ariaLabel?: string;
canCopy?: boolean;
text: string;
}
let { ariaLabel = 'Copy to clipboard', canCopy = true, text }: Props = $props();
export let ariaLabel: string = 'Copy to clipboard';
export let canCopy: boolean = true;
export let text: string;
</script>
<Copy
class="h-3 w-3 flex-shrink-0 cursor-{canCopy ? 'pointer' : 'not-allowed'}"
aria-label={ariaLabel}
<ActionIcon
icon={Copy}
tooltip={ariaLabel}
iconSize="h-4 w-4"
disabled={!canCopy}
onclick={() => canCopy && copyToClipboard(text)}
/>
@@ -1,27 +0,0 @@
<script lang="ts">
import { X } from '@lucide/svelte';
import { Button } from '$lib/components/ui/button';
interface Props {
id: string;
onRemove?: (id: string) => void;
class?: string;
iconSize?: number;
}
let { id, onRemove, class: className = '', iconSize = 3 }: Props = $props();
</script>
<Button
type="button"
variant="ghost"
size="icon-sm"
class="bg-white/20 p-0 hover:bg-white/30 {className}"
onclick={(e: MouseEvent) => {
e.stopPropagation();
onRemove?.(id);
}}
aria-label="Remove file"
>
<X class="h-{iconSize} w-{iconSize}" />
</Button>
@@ -1,46 +0,0 @@
<script lang="ts">
import { Eye } from '@lucide/svelte';
import { ActionIconCopyToClipboard } from '$lib/components/app';
import { FileTypeText } from '$lib/enums';
interface Props {
code: string;
language: string;
disabled?: boolean;
onPreview?: (code: string, language: string) => void;
}
let { code, language, disabled = false, onPreview }: Props = $props();
const showPreview = $derived(language?.toLowerCase() === FileTypeText.HTML);
function handlePreview() {
if (disabled) return;
onPreview?.(code, language);
}
</script>
<div class="code-block-actions">
<div class="copy-code-btn" class:opacity-50={disabled} class:!cursor-not-allowed={disabled}>
<ActionIconCopyToClipboard
text={code}
canCopy={!disabled}
ariaLabel={disabled ? 'Code incomplete' : 'Copy code'}
/>
</div>
{#if showPreview}
<button
class="preview-code-btn"
class:opacity-50={disabled}
class:!cursor-not-allowed={disabled}
title={disabled ? 'Code incomplete' : 'Preview code'}
aria-label="Preview code"
aria-disabled={disabled}
type="button"
onclick={handlePreview}
>
<Eye size={16} />
</button>
{/if}
</div>
@@ -9,11 +9,5 @@
/** Styled icon button for action triggers with tooltip. */
export { default as ActionIcon } from './ActionIcon.svelte';
/** Code block actions component (copy, preview). */
export { default as ActionIconsCodeBlock } from './ActionIconsCodeBlock.svelte';
/** Copy-to-clipboard icon button with click handler. */
/** Copy-to-clipboard icon button with clipboard logic. */
export { default as ActionIconCopyToClipboard } from './ActionIconCopyToClipboard.svelte';
/** Remove/delete icon button with X icon. */
export { default as ActionIconRemove } from './ActionIconRemove.svelte';
@@ -1,5 +1,4 @@
<script lang="ts">
import { cn } from '$lib/components/ui/utils';
import type { Snippet } from 'svelte';
interface Props {
@@ -13,10 +12,10 @@
</script>
<button
class={cn(
class={[
'inline-flex cursor-pointer items-center gap-1 rounded-sm bg-muted-foreground/15 px-1.5 py-0.75',
className
)}
]}
{onclick}
>
{#if icon}
@@ -1,39 +0,0 @@
<script lang="ts">
import { ModelModality } from '$lib/enums';
import { MODALITY_ICONS, MODALITY_LABELS } from '$lib/constants';
import { cn } from '$lib/components/ui/utils';
type DisplayableModality = ModelModality.VISION | ModelModality.AUDIO;
interface Props {
modalities: ModelModality[];
class?: string;
}
let { modalities, class: className = '' }: Props = $props();
// Filter to only modalities that have icons (VISION, AUDIO)
const displayableModalities = $derived(
modalities.filter(
(m): m is DisplayableModality => m === ModelModality.VISION || m === ModelModality.AUDIO
)
);
</script>
{#each displayableModalities as modality, index (index)}
{@const IconComponent = MODALITY_ICONS[modality]}
{@const label = MODALITY_LABELS[modality]}
<span
class={cn(
'inline-flex items-center gap-1 rounded-md bg-muted px-2 py-1 text-xs font-medium',
className
)}
>
{#if IconComponent}
<IconComponent class="h-3 w-3" />
{/if}
{label}
</span>
{/each}
@@ -0,0 +1,32 @@
<script lang="ts">
import { Eye, Mic } from '@lucide/svelte';
import { ModelModality } from '$lib/enums';
interface Props {
modalities: ModelModality[];
class?: string;
}
let { modalities, class: className = '' }: Props = $props();
</script>
{#each modalities as modality (modality)}
{#if modality === ModelModality.VISION || modality === ModelModality.AUDIO}
<span
class={[
'inline-flex items-center gap-1 rounded-md bg-muted px-2 py-1 text-xs font-medium',
className
]}
>
{#if modality === ModelModality.VISION}
<Eye class="h-3 w-3" />
Vision
{:else}
<Mic class="h-3 w-3" />
Audio
{/if}
</span>
{/if}
{/each}
@@ -6,11 +6,8 @@
*
*/
/** Badge displaying chat statistics (tokens, timing). */
export { default as BadgeChatStatistic } from './BadgeChatStatistic.svelte';
/** Generic info badge with optional tooltip and click handler. */
export { default as BadgeInfo } from './BadgeInfo.svelte';
/** Badge indicating model modality (vision, audio, tools). */
export { default as BadgeModality } from './BadgeModality.svelte';
export { default as BadgesModality } from './BadgesModality.svelte';
@@ -1,284 +0,0 @@
<script lang="ts">
import { Button } from '$lib/components/ui/button';
import * as Alert from '$lib/components/ui/alert';
import { SyntaxHighlightedCode } from '$lib/components/app';
import { FileText, Image, Music, FileIcon, Eye, Info } from '@lucide/svelte';
import {
isTextFile,
isImageFile,
isPdfFile,
isAudioFile,
getLanguageFromFilename,
createBase64DataUrl
} from '$lib/utils';
import { convertPDFToImage } from '$lib/utils/browser-only';
import { modelsStore } from '$lib/stores/models.svelte';
interface Props {
// Either an uploaded file or a stored attachment
uploadedFile?: ChatUploadedFile;
attachment?: DatabaseMessageExtra;
// For uploaded files
preview?: string;
name?: string;
textContent?: string;
// For checking vision modality
activeModelId?: string;
}
let { uploadedFile, attachment, preview, name, textContent, activeModelId }: Props = $props();
let hasVisionModality = $derived(
activeModelId ? modelsStore.modelSupportsVision(activeModelId) : false
);
let displayName = $derived(uploadedFile?.name || attachment?.name || name || 'Unknown File');
// Determine file type from uploaded file or attachment
let isAudio = $derived(isAudioFile(attachment, uploadedFile));
let isImage = $derived(isImageFile(attachment, uploadedFile));
let isPdf = $derived(isPdfFile(attachment, uploadedFile));
let isText = $derived(isTextFile(attachment, uploadedFile));
let displayPreview = $derived(
uploadedFile?.preview ||
(isImage && attachment && 'base64Url' in attachment ? attachment.base64Url : preview)
);
let displayTextContent = $derived(
uploadedFile?.textContent ||
(attachment && 'content' in attachment ? attachment.content : textContent)
);
let language = $derived(getLanguageFromFilename(displayName));
let IconComponent = $derived(() => {
if (isImage) return Image;
if (isText || isPdf) return FileText;
if (isAudio) return Music;
return FileIcon;
});
let pdfViewMode = $state<'text' | 'pages'>('pages');
let pdfImages = $state<string[]>([]);
let pdfImagesLoading = $state(false);
let pdfImagesError = $state<string | null>(null);
async function loadPdfImages() {
if (!isPdf || pdfImages.length > 0 || pdfImagesLoading) return;
pdfImagesLoading = true;
pdfImagesError = null;
try {
let file: File | null = null;
if (uploadedFile?.file) {
file = uploadedFile.file;
} else if (isPdf && attachment) {
// Check if we have pre-processed images
if (
'images' in attachment &&
attachment.images &&
Array.isArray(attachment.images) &&
attachment.images.length > 0
) {
pdfImages = attachment.images;
return;
}
// Convert base64 back to File for processing
if ('base64Data' in attachment && attachment.base64Data) {
const base64Data = attachment.base64Data;
const byteCharacters = atob(base64Data);
const byteNumbers = new Array(byteCharacters.length);
for (let i = 0; i < byteCharacters.length; i++) {
byteNumbers[i] = byteCharacters.charCodeAt(i);
}
const byteArray = new Uint8Array(byteNumbers);
file = new File([byteArray], displayName, { type: 'application/pdf' });
}
}
if (file) {
pdfImages = await convertPDFToImage(file);
} else {
throw new Error('No PDF file available for conversion');
}
} catch (error) {
pdfImagesError = error instanceof Error ? error.message : 'Failed to load PDF images';
} finally {
pdfImagesLoading = false;
}
}
export function reset() {
pdfImages = [];
pdfImagesLoading = false;
pdfImagesError = null;
pdfViewMode = 'pages';
}
$effect(() => {
if (isPdf && pdfViewMode === 'pages') {
loadPdfImages();
}
});
</script>
<div class="space-y-4">
<div class="flex items-center justify-end gap-6">
{#if isPdf}
<div class="flex items-center gap-2">
<Button
variant={pdfViewMode === 'text' ? 'default' : 'outline'}
size="sm"
onclick={() => (pdfViewMode = 'text')}
disabled={pdfImagesLoading}
>
<FileText class="mr-1 h-4 w-4" />
Text
</Button>
<Button
variant={pdfViewMode === 'pages' ? 'default' : 'outline'}
size="sm"
onclick={() => {
pdfViewMode = 'pages';
loadPdfImages();
}}
disabled={pdfImagesLoading}
>
{#if pdfImagesLoading}
<div
class="mr-1 h-4 w-4 animate-spin rounded-full border-2 border-current border-t-transparent"
></div>
{:else}
<Eye class="mr-1 h-4 w-4" />
{/if}
Pages
</Button>
</div>
{/if}
</div>
<div class="flex-1 overflow-auto">
{#if isImage && displayPreview}
<div class="flex items-center justify-center">
<img
src={displayPreview}
alt={displayName}
class="max-h-full rounded-lg object-contain shadow-lg"
/>
</div>
{:else if isPdf && pdfViewMode === 'pages'}
{#if !hasVisionModality && activeModelId}
<Alert.Root class="mb-4">
<Info class="h-4 w-4" />
<Alert.Title>Preview only</Alert.Title>
<Alert.Description>
<span class="inline-flex">
The selected model does not support vision. Only the extracted
<!-- svelte-ignore a11y_click_events_have_key_events -->
<!-- svelte-ignore a11y_no_static_element_interactions -->
<span class="mx-1 cursor-pointer underline" onclick={() => (pdfViewMode = 'text')}>
text
</span>
will be sent to the model.
</span>
</Alert.Description>
</Alert.Root>
{/if}
{#if pdfImagesLoading}
<div class="flex items-center justify-center p-8">
<div class="text-center">
<div
class="mx-auto mb-4 h-8 w-8 animate-spin rounded-full border-4 border-primary border-t-transparent"
></div>
<p class="text-muted-foreground">Converting PDF to images...</p>
</div>
</div>
{:else if pdfImagesError}
<div class="flex items-center justify-center p-8">
<div class="text-center">
<FileText class="mx-auto mb-4 h-16 w-16 text-muted-foreground" />
<p class="mb-4 text-muted-foreground">Failed to load PDF images</p>
<p class="text-sm text-muted-foreground">{pdfImagesError}</p>
<Button class="mt-4" onclick={() => (pdfViewMode = 'text')}>View as Text</Button>
</div>
</div>
{:else if pdfImages.length > 0}
<div class="max-h-[70vh] space-y-4 overflow-auto">
{#each pdfImages as image, index (image)}
<div class="text-center">
<p class="mb-2 text-sm text-muted-foreground">Page {index + 1}</p>
<img
src={image}
alt="PDF Page {index + 1}"
class="mx-auto max-w-full rounded-lg shadow-lg"
/>
</div>
{/each}
</div>
{:else}
<div class="flex items-center justify-center p-8">
<div class="text-center">
<FileText class="mx-auto mb-4 h-16 w-16 text-muted-foreground" />
<p class="mb-4 text-muted-foreground">No PDF pages available</p>
</div>
</div>
{/if}
{:else if (isText || (isPdf && pdfViewMode === 'text')) && displayTextContent}
<SyntaxHighlightedCode code={displayTextContent} {language} maxWidth="calc(69rem - 2rem)" />
{:else if isAudio}
<div class="flex items-center justify-center p-8">
<div class="w-full max-w-md text-center">
<Music class="mx-auto mb-4 h-16 w-16 text-muted-foreground" />
{#if uploadedFile?.preview}
<audio controls class="mb-4 w-full" src={uploadedFile.preview}>
Your browser does not support the audio element.
</audio>
{:else if isAudio && attachment && 'mimeType' in attachment && 'base64Data' in attachment}
<audio
controls
class="mb-4 w-full"
src={createBase64DataUrl(attachment.mimeType, attachment.base64Data)}
>
Your browser does not support the audio element.
</audio>
{:else}
<p class="mb-4 text-muted-foreground">Audio preview not available</p>
{/if}
<p class="text-sm text-muted-foreground">
{displayName}
</p>
</div>
</div>
{:else}
<div class="flex items-center justify-center p-8">
<div class="text-center">
{#if IconComponent}
<IconComponent class="mx-auto mb-4 h-16 w-16 text-muted-foreground" />
{/if}
<p class="mb-4 text-muted-foreground">Preview not available for this file type</p>
</div>
</div>
{/if}
</div>
</div>
@@ -1,165 +0,0 @@
<script lang="ts">
import { ActionIconRemove } from '$lib/components/app';
import { formatFileSize, getFileTypeLabel, getPreviewText, isTextFile } from '$lib/utils';
import { AttachmentType } from '$lib/enums';
interface Props {
class?: string;
id: string;
onClick?: (event?: MouseEvent) => void;
onRemove?: (id: string) => void;
name: string;
readonly?: boolean;
size?: number;
textContent?: string;
// Either uploaded file or stored attachment
uploadedFile?: ChatUploadedFile;
attachment?: DatabaseMessageExtra;
}
let {
class: className = '',
id,
onClick,
onRemove,
name,
readonly = false,
size,
textContent,
uploadedFile,
attachment
}: Props = $props();
let isText = $derived(isTextFile(attachment, uploadedFile));
let fileTypeLabel = $derived.by(() => {
if (uploadedFile?.type) {
return getFileTypeLabel(uploadedFile.type);
}
if (attachment) {
if ('mimeType' in attachment && attachment.mimeType) {
return getFileTypeLabel(attachment.mimeType);
}
if (attachment.type) {
return getFileTypeLabel(attachment.type);
}
}
return getFileTypeLabel(name);
});
let pdfProcessingMode = $derived.by(() => {
if (attachment?.type === AttachmentType.PDF) {
const pdfAttachment = attachment as DatabaseMessageExtraPdfFile;
return pdfAttachment.processedAsImages ? 'Sent as Image' : 'Sent as Text';
}
return null;
});
</script>
{#if isText}
{#if readonly}
<!-- Readonly mode (ChatMessage) -->
<button
class="cursor-pointer rounded-lg border border-border bg-muted p-3 transition-shadow hover:shadow-md {className} w-full max-w-2xl"
onclick={onClick}
aria-label={`Preview ${name}`}
type="button"
>
<div class="flex items-start gap-3">
<div class="flex min-w-0 flex-1 flex-col items-start text-left">
<span class="w-full truncate text-sm font-medium text-foreground">{name}</span>
{#if size}
<span class="text-xs text-muted-foreground">{formatFileSize(size)}</span>
{/if}
{#if textContent}
<div class="relative mt-2 w-full">
<div
class="overflow-hidden font-mono text-xs leading-relaxed break-words whitespace-pre-wrap text-muted-foreground"
>
{getPreviewText(textContent)}
</div>
{#if textContent.length > 150}
<div
class="pointer-events-none absolute right-0 bottom-0 left-0 h-6 bg-gradient-to-t from-muted to-transparent"
></div>
{/if}
</div>
{/if}
</div>
</div>
</button>
{:else}
<!-- Non-readonly mode (ChatForm) -->
<button
class="group relative rounded-lg border border-border bg-muted p-3 {className} {textContent
? 'max-h-24 max-w-72'
: 'max-w-36'} cursor-pointer text-left"
onclick={onClick}
>
<div class="absolute top-2 right-2 opacity-0 transition-opacity group-hover:opacity-100">
<ActionIconRemove {id} {onRemove} />
</div>
<div class="pr-8">
<span class="mb-3 block truncate text-sm font-medium text-foreground">{name}</span>
{#if textContent}
<div class="relative">
<div
class="overflow-hidden font-mono text-xs leading-relaxed break-words whitespace-pre-wrap text-muted-foreground"
style="max-height: 3rem; line-height: 1.2em;"
>
{getPreviewText(textContent)}
</div>
{#if textContent.length > 150}
<div
class="pointer-events-none absolute right-0 bottom-0 left-0 h-4 bg-gradient-to-t from-muted to-transparent"
></div>
{/if}
</div>
{/if}
</div>
</button>
{/if}
{:else}
<button
class="group flex items-center gap-3 rounded-lg border border-border bg-muted p-3 {className} relative"
onclick={onClick}
>
<div
class="flex h-8 w-8 items-center justify-center rounded bg-primary/10 text-xs font-medium text-primary"
>
{fileTypeLabel}
</div>
<div class="flex flex-col gap-0.5">
<span
class="max-w-24 truncate text-sm font-medium text-foreground {readonly
? ''
: 'group-hover:pr-6'} md:max-w-32"
>
{name}
</span>
{#if pdfProcessingMode}
<span class="text-left text-xs text-muted-foreground">{pdfProcessingMode}</span>
{:else if size}
<span class="text-left text-xs text-muted-foreground">{formatFileSize(size)}</span>
{/if}
</div>
{#if !readonly}
<div class="absolute top-2 right-2 opacity-0 transition-opacity group-hover:opacity-100">
<ActionIconRemove {id} {onRemove} />
</div>
{/if}
</button>
{/if}
@@ -1,287 +0,0 @@
<script lang="ts">
import {
ChatAttachmentMcpPrompt,
ChatAttachmentMcpResource,
ChatAttachmentThumbnailImage,
ChatAttachmentThumbnailFile,
HorizontalScrollCarousel,
DialogChatAttachmentPreview,
DialogChatAttachmentsViewAll,
DialogMcpResourcePreview
} from '$lib/components/app';
import { Button } from '$lib/components/ui/button';
import { AttachmentType } from '$lib/enums';
import type {
DatabaseMessageExtraMcpPrompt,
DatabaseMessageExtraMcpResource,
MCPResourceAttachment
} from '$lib/types';
import { getAttachmentDisplayItems } from '$lib/utils';
interface Props {
class?: string;
style?: string;
// For ChatMessage - stored attachments
attachments?: DatabaseMessageExtra[];
readonly?: boolean;
// For ChatForm - pending uploads
onFileRemove?: (fileId: string) => void;
uploadedFiles?: ChatUploadedFile[];
// Image size customization
imageClass?: string;
imageHeight?: string;
imageWidth?: string;
// Limit display to single row with "+ X more" button
limitToSingleRow?: boolean;
// For vision modality check
activeModelId?: string;
}
let {
class: className = '',
style = '',
attachments = [],
readonly = false,
onFileRemove,
uploadedFiles = $bindable([]),
// Default to small size for form previews
imageClass = '',
imageHeight = 'h-24',
imageWidth = 'w-auto',
limitToSingleRow = false,
activeModelId
}: Props = $props();
let displayItems = $derived(getAttachmentDisplayItems({ uploadedFiles, attachments }));
let carouselRef: HorizontalScrollCarousel | undefined = $state();
let isScrollable = $state(false);
let previewDialogOpen = $state(false);
let previewItem = $state<ChatAttachmentPreviewItem | null>(null);
let mcpResourcePreviewOpen = $state(false);
let mcpResourcePreviewExtra = $state<DatabaseMessageExtraMcpResource | null>(null);
let showViewAll = $derived(limitToSingleRow && displayItems.length > 0 && isScrollable);
let viewAllDialogOpen = $state(false);
function openPreview(item: ChatAttachmentDisplayItem, event?: MouseEvent) {
event?.stopPropagation();
event?.preventDefault();
previewItem = {
uploadedFile: item.uploadedFile,
attachment: item.attachment,
preview: item.preview,
name: item.name,
size: item.size,
textContent: item.textContent
};
previewDialogOpen = true;
}
function openMcpResourcePreview(extra: DatabaseMessageExtraMcpResource) {
mcpResourcePreviewExtra = extra;
mcpResourcePreviewOpen = true;
}
function toMcpResourceAttachment(
extra: DatabaseMessageExtraMcpResource,
id: string
): MCPResourceAttachment {
return {
id,
resource: {
uri: extra.uri,
name: extra.name,
title: extra.name,
serverName: extra.serverName
}
};
}
$effect(() => {
if (carouselRef && displayItems.length) {
carouselRef.resetScroll();
}
});
</script>
{#if displayItems.length > 0}
<div class={className} {style}>
{#if limitToSingleRow}
<HorizontalScrollCarousel
bind:this={carouselRef}
onScrollableChange={(scrollable) => (isScrollable = scrollable)}
>
{#each displayItems as item (item.id)}
{#if item.isMcpPrompt}
{@const mcpPrompt =
item.attachment?.type === AttachmentType.MCP_PROMPT
? (item.attachment as DatabaseMessageExtraMcpPrompt)
: item.uploadedFile?.mcpPrompt
? {
type: AttachmentType.MCP_PROMPT as const,
name: item.name,
serverName: item.uploadedFile.mcpPrompt.serverName,
promptName: item.uploadedFile.mcpPrompt.promptName,
content: item.textContent ?? '',
arguments: item.uploadedFile.mcpPrompt.arguments
}
: null}
{#if mcpPrompt}
<ChatAttachmentMcpPrompt
class="max-w-[300px] min-w-[200px] flex-shrink-0 {limitToSingleRow
? 'first:ml-4 last:mr-4'
: ''}"
prompt={mcpPrompt}
{readonly}
isLoading={item.isLoading}
loadError={item.loadError}
onRemove={onFileRemove ? () => onFileRemove(item.id) : undefined}
/>
{/if}
{:else if item.isMcpResource && item.attachment?.type === AttachmentType.MCP_RESOURCE}
{@const mcpResource = item.attachment as DatabaseMessageExtraMcpResource}
<ChatAttachmentMcpResource
class="flex-shrink-0 {limitToSingleRow ? 'first:ml-4 last:mr-4' : ''}"
attachment={toMcpResourceAttachment(mcpResource, item.id)}
onClick={() => openMcpResourcePreview(mcpResource)}
/>
{:else if item.isImage && item.preview}
<ChatAttachmentThumbnailImage
class="flex-shrink-0 cursor-pointer {limitToSingleRow ? 'first:ml-4 last:mr-4' : ''}"
id={item.id}
name={item.name}
preview={item.preview}
{readonly}
onRemove={onFileRemove}
height={imageHeight}
width={imageWidth}
{imageClass}
onClick={(event) => openPreview(item, event)}
/>
{:else}
<ChatAttachmentThumbnailFile
class="flex-shrink-0 cursor-pointer {limitToSingleRow ? 'first:ml-4 last:mr-4' : ''}"
id={item.id}
name={item.name}
size={item.size}
{readonly}
onRemove={onFileRemove}
textContent={item.textContent}
attachment={item.attachment}
uploadedFile={item.uploadedFile}
onClick={(event) => openPreview(item, event)}
/>
{/if}
{/each}
</HorizontalScrollCarousel>
{#if showViewAll}
<div class="mt-2 -mr-2 flex justify-end px-4">
<Button
type="button"
variant="ghost"
size="sm"
class="h-6 text-xs text-muted-foreground hover:text-foreground"
onclick={() => (viewAllDialogOpen = true)}
>
View all ({displayItems.length})
</Button>
</div>
{/if}
{:else}
<div class="flex flex-wrap items-start justify-end gap-3">
{#each displayItems as item (item.id)}
{#if item.isMcpPrompt}
{@const mcpPrompt =
item.attachment?.type === AttachmentType.MCP_PROMPT
? (item.attachment as DatabaseMessageExtraMcpPrompt)
: item.uploadedFile?.mcpPrompt
? {
type: AttachmentType.MCP_PROMPT as const,
name: item.name,
serverName: item.uploadedFile.mcpPrompt.serverName,
promptName: item.uploadedFile.mcpPrompt.promptName,
content: item.textContent ?? '',
arguments: item.uploadedFile.mcpPrompt.arguments
}
: null}
{#if mcpPrompt}
<ChatAttachmentMcpPrompt
class="max-w-[300px] min-w-[200px]"
prompt={mcpPrompt}
{readonly}
isLoading={item.isLoading}
loadError={item.loadError}
onRemove={onFileRemove ? () => onFileRemove(item.id) : undefined}
/>
{/if}
{:else if item.isMcpResource && item.attachment?.type === AttachmentType.MCP_RESOURCE}
{@const mcpResource = item.attachment as DatabaseMessageExtraMcpResource}
<ChatAttachmentMcpResource
attachment={toMcpResourceAttachment(mcpResource, item.id)}
onClick={() => openMcpResourcePreview(mcpResource)}
/>
{:else if item.isImage && item.preview}
<ChatAttachmentThumbnailImage
class="cursor-pointer"
id={item.id}
name={item.name}
preview={item.preview}
{readonly}
onRemove={onFileRemove}
height={imageHeight}
width={imageWidth}
{imageClass}
onClick={(event) => openPreview(item, event)}
/>
{:else}
<ChatAttachmentThumbnailFile
class="cursor-pointer"
id={item.id}
name={item.name}
size={item.size}
{readonly}
onRemove={onFileRemove}
textContent={item.textContent}
attachment={item.attachment}
uploadedFile={item.uploadedFile}
onClick={(event?: MouseEvent) => openPreview(item, event)}
/>
{/if}
{/each}
</div>
{/if}
</div>
{/if}
{#if previewItem}
<DialogChatAttachmentPreview
bind:open={previewDialogOpen}
uploadedFile={previewItem.uploadedFile}
attachment={previewItem.attachment}
preview={previewItem.preview}
name={previewItem.name}
size={previewItem.size}
textContent={previewItem.textContent}
{activeModelId}
/>
{/if}
<DialogChatAttachmentsViewAll
bind:open={viewAllDialogOpen}
{uploadedFiles}
{attachments}
{readonly}
{onFileRemove}
imageHeight="h-64"
{imageClass}
{activeModelId}
/>
{#if mcpResourcePreviewExtra}
<DialogMcpResourcePreview bind:open={mcpResourcePreviewOpen} extra={mcpResourcePreviewExtra} />
{/if}
@@ -0,0 +1,119 @@
<script lang="ts">
import {
ChatAttachmentsListItem,
DialogChatAttachmentsPreview,
DialogMcpResourcePreview,
HorizontalScrollCarousel
} from '$lib/components/app';
import type { DatabaseMessageExtraMcpResource } from '$lib/types';
import { getAttachmentDisplayItems, isMcpPrompt, isMcpResource } from '$lib/utils';
interface Props {
class?: string;
style?: string;
// For ChatMessage - stored attachments
attachments?: DatabaseMessageExtra[];
readonly?: boolean;
// For ChatForm - pending uploads
onFileRemove?: (fileId: string) => void;
uploadedFiles?: ChatUploadedFile[];
// Image size customization
imageClass?: string;
imageHeight?: string;
imageWidth?: string;
// Limit display to single row with "+ X more" button
limitToSingleRow?: boolean;
// For vision modality check
activeModelId?: string;
}
let {
class: className = '',
style = '',
attachments = [],
readonly = false,
onFileRemove,
uploadedFiles = $bindable([]),
// Default to small size for form previews
imageClass = '',
imageHeight = 'h-24',
imageWidth = 'w-auto',
limitToSingleRow = false,
activeModelId
}: Props = $props();
let carouselRef: HorizontalScrollCarousel | undefined = $state();
let mcpResourcePreviewOpen = $state(false);
let mcpResourcePreviewExtra = $state<DatabaseMessageExtraMcpResource | null>(null);
let previewFocusIndex = $state(0);
let viewAllDialogOpen = $state(false);
let displayItems = $derived(getAttachmentDisplayItems({ uploadedFiles, attachments }));
function openPreview(item: ChatAttachmentDisplayItem, event?: MouseEvent) {
event?.stopPropagation();
event?.preventDefault();
// Find the index of the clicked item among non-MCP attachments
const nonMcpItems = displayItems.filter((i) => !isMcpPrompt(i) && !isMcpResource(i));
const index = nonMcpItems.findIndex((i) => i.id === item.id);
previewFocusIndex = index >= 0 ? index : 0;
viewAllDialogOpen = true;
}
function openMcpResourcePreview(extra: DatabaseMessageExtraMcpResource) {
mcpResourcePreviewExtra = extra;
mcpResourcePreviewOpen = true;
}
$effect(() => {
if (carouselRef && displayItems.length) {
carouselRef.resetScroll();
}
});
</script>
{#snippet attachmentitem(item: ChatAttachmentDisplayItem)}
<ChatAttachmentsListItem
{imageClass}
{imageHeight}
{imageWidth}
{item}
{limitToSingleRow}
{onFileRemove}
onMcpResourcePreview={openMcpResourcePreview}
onPreview={(i: ChatAttachmentDisplayItem, event?: MouseEvent) => openPreview(i, event)}
{readonly}
/>
{/snippet}
{#if displayItems.length > 0}
<div class={className} {style}>
{#if limitToSingleRow}
<HorizontalScrollCarousel bind:this={carouselRef}>
{#each displayItems as item (item.id)}
{@render attachmentitem(item)}
{/each}
</HorizontalScrollCarousel>
{:else}
<div class="flex flex-wrap items-start justify-end gap-3">
{#each displayItems as item (item.id)}
{@render attachmentitem(item)}
{/each}
</div>
{/if}
</div>
{/if}
<DialogChatAttachmentsPreview
{activeModelId}
{attachments}
bind:open={viewAllDialogOpen}
{previewFocusIndex}
{uploadedFiles}
/>
{#if mcpResourcePreviewExtra}
<DialogMcpResourcePreview extra={mcpResourcePreviewExtra} bind:open={mcpResourcePreviewOpen} />
{/if}
@@ -0,0 +1,132 @@
<script lang="ts">
import {
ChatAttachmentsListItemMcpPrompt,
ChatAttachmentsListItemMcpResource,
ChatAttachmentsListItemThumbnailImage,
ChatAttachmentsListItemThumbnailFile
} from '$lib/components/app';
import { AttachmentType } from '$lib/enums';
import type {
ChatAttachmentDisplayItem,
DatabaseMessageExtraMcpPrompt,
DatabaseMessageExtraMcpResource,
MCPResourceAttachment
} from '$lib/types';
import { isMcpPrompt, isMcpResource, isPdfFile } from '$lib/utils';
interface Props {
class?: string;
imageClass?: string;
imageHeight?: string;
imageWidth?: string;
item: ChatAttachmentDisplayItem;
limitToSingleRow?: boolean;
onFileRemove?: (fileId: string) => void;
onMcpResourcePreview?: (extra: DatabaseMessageExtraMcpResource) => void;
onPreview?: (item: ChatAttachmentDisplayItem) => void;
readonly?: boolean;
}
let {
class: className = '',
imageClass = '',
imageHeight = 'h-24',
imageWidth = 'w-auto',
item,
limitToSingleRow = false,
onFileRemove,
onMcpResourcePreview,
onPreview,
readonly = false
}: Props = $props();
const scrollClasses = $derived(limitToSingleRow ? 'first:ml-4 last:mr-4' : '');
function toMcpResourceAttachment(
extra: DatabaseMessageExtraMcpResource,
id: string
): MCPResourceAttachment {
return {
id,
resource: {
uri: extra.uri,
name: extra.name,
title: extra.name,
serverName: extra.serverName
}
};
}
</script>
{#if isMcpPrompt(item)}
{@const mcpPrompt =
item.attachment?.type === AttachmentType.MCP_PROMPT
? (item.attachment as DatabaseMessageExtraMcpPrompt)
: item.uploadedFile?.mcpPrompt
? {
type: AttachmentType.MCP_PROMPT as const,
name: item.name,
serverName: item.uploadedFile.mcpPrompt.serverName,
promptName: item.uploadedFile.mcpPrompt.promptName,
content: item.textContent ?? '',
arguments: item.uploadedFile.mcpPrompt.arguments
}
: null}
{#if mcpPrompt}
<ChatAttachmentsListItemMcpPrompt
class="max-w-[300px] min-w-[200px] flex-shrink-0 {className} {scrollClasses}"
prompt={mcpPrompt}
{readonly}
isLoading={item.isLoading}
loadError={item.loadError}
onRemove={onFileRemove ? () => onFileRemove(item.id) : undefined}
/>
{/if}
{:else if isMcpResource(item)}
{@const mcpResource = item.attachment as DatabaseMessageExtraMcpResource}
<ChatAttachmentsListItemMcpResource
class="flex-shrink-0 {className} {scrollClasses}"
attachment={toMcpResourceAttachment(mcpResource, item.id)}
onclick={() => onMcpResourcePreview?.(mcpResource)}
/>
{:else if item.isImage && item.preview}
<ChatAttachmentsListItemThumbnailImage
class="flex-shrink-0 cursor-pointer {className} {scrollClasses}"
id={item.id}
name={item.name}
preview={item.preview}
{readonly}
onRemove={onFileRemove}
height={imageHeight}
width={imageWidth}
{imageClass}
onclick={() => onPreview?.(item)}
/>
{:else if isPdfFile(item.attachment, item.uploadedFile)}
<ChatAttachmentsListItemThumbnailFile
class="flex-shrink-0 cursor-pointer {className} {scrollClasses}"
id={item.id}
name={item.name}
size={item.size}
{readonly}
onRemove={onFileRemove}
textContent={item.textContent}
attachment={item.attachment}
uploadedFile={item.uploadedFile}
onclick={() => onPreview?.(item)}
/>
{:else}
<ChatAttachmentsListItemThumbnailFile
class="flex-shrink-0 cursor-pointer {className} {scrollClasses}"
id={item.id}
name={item.name}
size={item.size}
{readonly}
onRemove={onFileRemove}
textContent={item.textContent}
attachment={item.attachment}
uploadedFile={item.uploadedFile}
onclick={() => onPreview?.(item)}
/>
{/if}
@@ -1,40 +1,41 @@
<script lang="ts">
import { ChatMessageMcpPromptContent, ActionIconRemove } from '$lib/components/app';
import { ChatMessageMcpPromptContent, ActionIcon } from '$lib/components/app';
import { X } from '@lucide/svelte';
import type { DatabaseMessageExtraMcpPrompt } from '$lib/types';
import { McpPromptVariant } from '$lib/enums';
interface Props {
class?: string;
prompt: DatabaseMessageExtraMcpPrompt;
readonly?: boolean;
isLoading?: boolean;
loadError?: string;
onRemove?: () => void;
prompt: DatabaseMessageExtraMcpPrompt;
readonly?: boolean;
}
let {
class: className = '',
prompt,
readonly = false,
isLoading = false,
loadError,
onRemove
onRemove,
prompt,
readonly = false
}: Props = $props();
</script>
<div class="group relative {className}">
<ChatMessageMcpPromptContent
{prompt}
variant={McpPromptVariant.ATTACHMENT}
{isLoading}
{loadError}
{prompt}
variant={McpPromptVariant.ATTACHMENT}
/>
{#if !readonly && onRemove}
<div
class="absolute top-10 right-2 flex items-center justify-center opacity-0 transition-opacity group-hover:opacity-100"
>
<ActionIconRemove id={prompt.name} onRemove={() => onRemove?.()} />
<ActionIcon icon={X} tooltip="Remove" stopPropagationOnClick onclick={() => onRemove?.()} />
</div>
{/if}
</div>
@@ -1,46 +1,47 @@
<script lang="ts">
import { Loader2, AlertCircle } from '@lucide/svelte';
import { cn } from '$lib/components/ui/utils';
import { mcpStore } from '$lib/stores/mcp.svelte';
import type { MCPResourceAttachment } from '$lib/types';
import * as Tooltip from '$lib/components/ui/tooltip';
import { ActionIconRemove } from '$lib/components/app';
import { ActionIcon } from '$lib/components/app';
import { X } from '@lucide/svelte';
import { getResourceIcon, getResourceDisplayName } from '$lib/utils';
interface Props {
attachment: MCPResourceAttachment;
onRemove?: (attachmentId: string) => void;
onClick?: () => void;
class?: string;
onclick?: () => void;
onRemove?: (attachmentId: string) => void;
}
let { attachment, onRemove, onClick, class: className }: Props = $props();
function getStatusClass(attachment: MCPResourceAttachment): string {
if (attachment.error) return 'border-red-500/50 bg-red-500/10';
if (attachment.loading) return 'border-border/50 bg-muted/30';
return 'border-border/50 bg-muted/30';
}
let { attachment, class: className, onclick, onRemove }: Props = $props();
const ResourceIcon = $derived(
getResourceIcon(attachment.resource.mimeType, attachment.resource.uri)
);
const serverName = $derived(mcpStore.getServerDisplayName(attachment.resource.serverName));
const favicon = $derived(mcpStore.getServerFavicon(attachment.resource.serverName));
function getStatusClass(attachment: MCPResourceAttachment): string {
if (attachment.error) return 'border-red-500/50 bg-red-500/10';
if (attachment.loading) return 'border-border/50 bg-muted/30';
return 'border-border/50 bg-muted/30';
}
</script>
<Tooltip.Root>
<Tooltip.Trigger>
<button
type="button"
class={cn(
class={[
'flex flex-shrink-0 items-center gap-1.5 rounded-md border px-2 py-0.75 text-sm transition-colors',
getStatusClass(attachment),
onClick && 'cursor-pointer hover:bg-muted/50',
onclick && 'cursor-pointer hover:bg-muted/50',
className
)}
onclick={onClick}
disabled={!onClick}
]}
disabled={!onclick}
{onclick}
type="button"
>
{#if attachment.loading}
<Loader2 class="h-3 w-3 animate-spin text-muted-foreground" />
@@ -55,11 +56,13 @@
</span>
{#if onRemove}
<ActionIconRemove
<ActionIcon
class="-my-2 -mr-1.5 bg-transparent"
iconSize={2}
id={attachment.id}
{onRemove}
icon={X}
iconSize="h-2 w-2"
onclick={() => onRemove?.(attachment.id)}
stopPropagationOnClick
tooltip="Remove"
/>
{/if}
</button>
@@ -69,12 +72,12 @@
<div class="flex items-center gap-1 text-xs">
{#if favicon}
<img
src={favicon}
alt=""
alt={attachment.resource.serverName}
class="h-3 w-3 shrink-0 rounded-sm"
onerror={(e) => {
(e.currentTarget as HTMLImageElement).style.display = 'none';
}}
src={favicon}
/>
{/if}
@@ -0,0 +1,174 @@
<script lang="ts">
import { X } from '@lucide/svelte';
import {
formatFileSize,
getFileTypeLabel,
getPreviewText,
isPdfFile,
isTextFile
} from '$lib/utils';
import { ActionIcon } from '$lib/components/app';
import { AttachmentType } from '$lib/enums';
interface Props {
attachment?: DatabaseMessageExtra;
class?: string;
id: string;
onclick?: (event: MouseEvent) => void;
onRemove?: (id: string) => void;
name: string;
readonly?: boolean;
size?: number;
textContent?: string;
// Either uploaded file or stored attachment
uploadedFile?: ChatUploadedFile;
}
let {
attachment,
class: className = '',
id,
onclick,
onRemove,
name,
readonly = false,
size,
textContent,
uploadedFile
}: Props = $props();
let isPdf = $derived(isPdfFile(attachment, uploadedFile));
let isPdfWithContent = $derived(isPdf && !!textContent);
let isText = $derived(isTextFile(attachment, uploadedFile));
let isTextWithContent = $derived(isText && !!textContent);
let fileTypeLabel = $derived.by(() => {
if (uploadedFile?.type) {
return getFileTypeLabel(uploadedFile.type);
}
if (attachment) {
if ('mimeType' in attachment && attachment.mimeType) {
return getFileTypeLabel(attachment.mimeType);
}
if (attachment.type) {
return getFileTypeLabel(attachment.type);
}
}
return getFileTypeLabel(name);
});
let pdfProcessingMode = $derived.by(() => {
if (attachment?.type === AttachmentType.PDF) {
const pdfAttachment = attachment as DatabaseMessageExtraPdfFile;
return pdfAttachment.processedAsImages ? 'Sent as Image' : 'Sent as Text';
}
return null;
});
</script>
{#snippet textPreview(content: string)}
<div class="relative">
<div
class="font-mono text-xs leading-relaxed break-words whitespace-pre-wrap text-muted-foreground {!readonly
? 'max-h-3rem line-height-1.2'
: ''}"
>
{getPreviewText(content)}
</div>
{#if content.length > 150}
<div
class="pointer-events-none absolute right-0 bottom-0 left-0 h-4 bg-gradient-to-t from-muted to-transparent {readonly
? 'h-6'
: ''}"
></div>
{/if}
</div>
{/snippet}
{#snippet removeButton()}
<div class="absolute top-2 right-2 opacity-0 transition-opacity group-hover:opacity-100">
<ActionIcon icon={X} tooltip="Remove" stopPropagationOnClick onclick={() => onRemove?.(id)} />
</div>
{/snippet}
{#snippet fileIcon()}
<div
class="flex h-8 w-8 items-center justify-center rounded bg-primary/10 text-xs font-medium text-primary"
>
{fileTypeLabel}
</div>
{/snippet}
{#snippet info(text: string | undefined)}
{#if text}
<span class="text-xs text-muted-foreground">{text}</span>
{/if}
{/snippet}
{#if isTextWithContent || isPdfWithContent}
<button
aria-label={readonly ? `Preview ${name}` : undefined}
class="rounded-lg border border-border bg-muted p-3 {className} cursor-pointer {readonly
? 'w-full max-w-2xl transition-shadow hover:shadow-md'
: `group relative text-left ${textContent ? 'max-h-24 max-w-72' : 'max-w-36'}`} overflow-hidden"
{onclick}
type="button"
>
{#if !readonly}
{@render removeButton()}
{/if}
<div class={[!readonly && 'pr-8', 'overflow-hidden']}>
{#if readonly}
<div class="flex items-start gap-3">
<div class="flex min-w-0 flex-1 flex-col items-start text-left">
<span class="w-full truncate text-sm font-medium text-foreground">{name}</span>
{@render info(pdfProcessingMode || (size ? formatFileSize(size) : undefined))}
{#if textContent}
{@render textPreview(textContent)}
{/if}
</div>
</div>
{:else}
<span class="mb-3 block truncate text-sm font-medium text-foreground">{name}</span>
{#if textContent}
{@render textPreview(textContent)}
{/if}
{/if}
</div>
</button>
{:else}
<button
class="group flex items-center gap-3 rounded-lg border border-border bg-muted p-3 {className} relative"
{onclick}
type="button"
>
{@render fileIcon()}
<div class="flex flex-col items-start gap-0.5">
<span
class="max-w-24 truncate text-sm font-medium text-foreground {readonly
? ''
: 'group-hover:pr-6'} md:max-w-32"
>
{name}
</span>
{@render info(pdfProcessingMode || (size ? formatFileSize(size) : undefined))}
</div>
{#if !readonly}
{@render removeButton()}
{/if}
</button>
{/if}
@@ -1,64 +1,65 @@
<script lang="ts">
import { ActionIconRemove } from '$lib/components/app';
import { ActionIcon } from '$lib/components/app';
import { X } from '@lucide/svelte';
interface Props {
class?: string;
height?: string;
id: string;
imageClass?: string;
onclick?: (event?: MouseEvent) => void;
onRemove?: (id: string) => void;
name: string;
preview: string;
readonly?: boolean;
onRemove?: (id: string) => void;
onClick?: (event?: MouseEvent) => void;
class?: string;
// Customizable size props
width?: string;
height?: string;
imageClass?: string;
}
let {
class: className = '',
height = 'h-16',
id,
imageClass = '',
onclick,
onRemove,
name,
preview,
readonly = false,
onRemove,
onClick,
class: className = '',
// Default to small size for form previews
width = 'w-auto',
height = 'h-16',
imageClass = ''
width = 'w-auto'
}: Props = $props();
</script>
{#snippet image()}
<img src={preview} alt={name} class="{height} {width} cursor-pointer object-cover {imageClass}" />
{/snippet}
<div
class="group relative overflow-hidden rounded-lg bg-muted shadow-lg dark:border dark:border-muted {className}"
>
{#if onClick}
{#if onclick}
<button
type="button"
class="block h-full w-full rounded-lg focus:ring-2 focus:ring-primary focus:ring-offset-2 focus:outline-none"
onclick={onClick}
aria-label="Preview {name}"
class="block h-full w-full rounded-lg focus:ring-2 focus:ring-primary focus:ring-offset-2 focus:outline-none"
{onclick}
type="button"
>
<img
src={preview}
alt={name}
class="{height} {width} cursor-pointer object-cover {imageClass}"
/>
{@render image()}
</button>
{:else}
<img
src={preview}
alt={name}
class="{height} {width} cursor-pointer object-cover {imageClass}"
/>
{@render image()}
{/if}
{#if !readonly}
<div
class="absolute top-1 right-1 flex items-center justify-center opacity-0 transition-opacity group-hover:opacity-100"
>
<ActionIconRemove {id} {onRemove} class="text-white" />
<ActionIcon
class="text-white"
icon={X}
onclick={() => onRemove?.(id)}
stopPropagationOnClick
tooltip="Remove"
/>
</div>
{/if}
</div>
@@ -0,0 +1,190 @@
<script lang="ts">
import {
ChatAttachmentsPreviewCurrentItem,
ChatAttachmentsPreviewFileInfo,
ChatAttachmentsPreviewNavButtons,
ChatAttachmentsPreviewThumbnailStrip
} from '$lib/components/app';
import { modelsStore } from '$lib/stores/models.svelte';
import {
createBase64DataUrl,
formatFileSize,
getAttachmentDisplayItems,
getLanguageFromFilename,
isAudioFile,
isImageFile,
isMcpPrompt,
isMcpResource,
isPdfFile,
isTextFile
} from '$lib/utils';
interface PreviewItem {
id: string;
name: string;
size?: number;
preview?: string;
uploadedFile?: ChatUploadedFile;
attachment?: DatabaseMessageExtra;
textContent?: string;
isImage: boolean;
isAudio: boolean;
}
interface Props {
uploadedFiles?: ChatUploadedFile[];
attachments?: DatabaseMessageExtra[];
activeModelId?: string;
class?: string;
previewFocusIndex?: number;
}
let {
uploadedFiles = [],
attachments = [],
activeModelId,
class: className = '',
previewFocusIndex = 0
}: Props = $props();
let allItems = $derived(
getAttachmentDisplayItems({ uploadedFiles, attachments })
.filter((item) => !isMcpPrompt(item) && !isMcpResource(item))
.map(
(item): PreviewItem => ({
...item,
isImage: isImageFile(item.attachment, item.uploadedFile),
isAudio: isAudioFile(item.attachment, item.uploadedFile)
})
)
);
let currentIndex = $state(0);
$effect(() => {
if (previewFocusIndex >= 0 && previewFocusIndex < allItems.length) {
currentIndex = previewFocusIndex;
}
});
$effect(() => {
const handler = (e: Event) => {
const delta = (e as CustomEvent).detail;
if (delta < 0) {
currentIndex = currentIndex > 0 ? currentIndex - 1 : allItems.length - 1;
} else {
currentIndex = currentIndex < allItems.length - 1 ? currentIndex + 1 : 0;
}
};
document.addEventListener('chat-attachments-nav', handler);
return () => document.removeEventListener('chat-attachments-nav', handler);
});
$effect(() => {
const index = currentIndex;
setTimeout(() => {
const thumbnail = document.querySelector(`[data-thumbnail-index="${index}"]`);
thumbnail?.scrollIntoView({ behavior: 'smooth', inline: 'center', block: 'nearest' });
}, 0);
});
let currentItem = $derived(allItems[currentIndex] ?? null);
let displayName = $derived(
currentItem?.name ||
currentItem?.uploadedFile?.name ||
currentItem?.attachment?.name ||
'Unknown File'
);
let isAudio = $derived(
currentItem ? isAudioFile(currentItem.attachment, currentItem.uploadedFile) : false
);
let isImage = $derived(
currentItem ? isImageFile(currentItem.attachment, currentItem.uploadedFile) : false
);
let isPdf = $derived(
currentItem ? isPdfFile(currentItem.attachment, currentItem.uploadedFile) : false
);
let isText = $derived(
currentItem ? isTextFile(currentItem.attachment, currentItem.uploadedFile) : false
);
let displayPreview = $derived(
currentItem?.uploadedFile?.preview ||
(isImage && currentItem?.attachment && 'base64Url' in currentItem.attachment
? currentItem.attachment.base64Url
: currentItem?.preview)
);
let displayTextContent = $derived(
currentItem?.uploadedFile?.textContent ||
(currentItem?.attachment && 'content' in currentItem.attachment
? currentItem.attachment.content
: currentItem?.textContent)
);
let language = $derived(getLanguageFromFilename(displayName));
let fileSize = $derived(currentItem?.size ? formatFileSize(currentItem.size) : '');
let hasVisionModality = $derived(
currentItem && activeModelId ? modelsStore.modelSupportsVision(activeModelId) : false
);
let audioSrc = $derived(
isAudio && currentItem
? (currentItem.uploadedFile?.preview ??
(currentItem.attachment &&
'mimeType' in currentItem.attachment &&
'base64Data' in currentItem.attachment
? createBase64DataUrl(
currentItem.attachment.mimeType,
currentItem.attachment.base64Data
)
: null))
: null
);
export function prev() {
currentIndex = currentIndex > 0 ? currentIndex - 1 : allItems.length - 1;
}
export function next() {
currentIndex = currentIndex < allItems.length - 1 ? currentIndex + 1 : 0;
}
function onNavigate(index: number) {
currentIndex = index;
}
</script>
<div class="{className} flex flex-col text-white">
<div class="relative flex min-h-0 flex-1 items-center justify-center overflow-hidden">
<ChatAttachmentsPreviewNavButtons onPrev={prev} onNext={next} show={allItems.length > 1} />
<div class="flex h-full w-full flex-col items-center justify-start overflow-auto py-4">
{#if currentItem}
<ChatAttachmentsPreviewFileInfo {displayName} {fileSize} />
<ChatAttachmentsPreviewCurrentItem
{currentItem}
{isImage}
{isAudio}
{isPdf}
{isText}
{displayPreview}
{displayTextContent}
{audioSrc}
{language}
{hasVisionModality}
{activeModelId}
/>
{/if}
<ChatAttachmentsPreviewThumbnailStrip items={allItems} {currentIndex} {onNavigate} />
</div>
</div>
</div>
@@ -0,0 +1,65 @@
<script lang="ts">
import type { ChatAttachmentDisplayItem } from '$lib/types';
import { Image, Music, FileText, FileIcon } from '@lucide/svelte';
import ChatAttachmentsPreviewCurrentItemPdf from './ChatAttachmentsPreviewCurrentItemPdf.svelte';
import ChatAttachmentsPreviewCurrentItemImage from './ChatAttachmentsPreviewCurrentItemImage.svelte';
import ChatAttachmentsPreviewCurrentItemAudio from './ChatAttachmentsPreviewCurrentItemAudio.svelte';
import ChatAttachmentsPreviewCurrentItemText from './ChatAttachmentsPreviewCurrentItemText.svelte';
import ChatAttachmentsPreviewCurrentItemUnavailable from './ChatAttachmentsPreviewCurrentItemUnavailable.svelte';
interface Props {
currentItem: ChatAttachmentDisplayItem | null;
isImage: boolean;
isAudio: boolean;
isPdf: boolean;
isText: boolean;
displayPreview: string | undefined;
displayTextContent: string | undefined;
audioSrc: string | null;
language: string;
hasVisionModality: boolean;
activeModelId?: string;
}
let {
currentItem,
isImage,
isAudio,
isPdf,
isText,
displayPreview,
displayTextContent,
audioSrc,
language,
hasVisionModality,
activeModelId
}: Props = $props();
let IconComponent = $derived(
isImage ? Image : isText || isPdf ? FileText : isAudio ? Music : FileIcon
);
let isUnavailable = $derived(!isPdf && !isImage && !(isText && displayTextContent) && !isAudio);
</script>
{#if currentItem}
{#key currentItem.id}
{#if isPdf}
<ChatAttachmentsPreviewCurrentItemPdf
{currentItem}
displayName={currentItem.name}
{displayTextContent}
{hasVisionModality}
{activeModelId}
/>
{:else if isImage}
<ChatAttachmentsPreviewCurrentItemImage {currentItem} {displayPreview} />
{:else if isText && displayTextContent}
<ChatAttachmentsPreviewCurrentItemText {displayTextContent} {language} />
{:else if isAudio}
<ChatAttachmentsPreviewCurrentItemAudio {currentItem} {audioSrc} />
{:else if isUnavailable}
<ChatAttachmentsPreviewCurrentItemUnavailable {IconComponent} />
{/if}
{/key}
{/if}
@@ -0,0 +1,26 @@
<script lang="ts">
import { Music } from '@lucide/svelte';
interface Props {
currentItem: { name?: string } | null;
audioSrc: string | null;
}
let { currentItem, audioSrc }: Props = $props();
</script>
<div class="flex flex-1 items-center justify-center p-8">
<div class="w-full max-w-md text-center">
<Music class="mx-auto mb-4 h-16 w-16 text-white/50" />
{#if audioSrc}
<audio controls class="mb-4 w-full" src={audioSrc}>
Your browser does not support the audio element.
</audio>
{:else}
<p class="mb-4 text-white/70">Audio preview not available</p>
{/if}
<p class="text-sm text-white/50">{currentItem?.name || 'Audio'}</p>
</div>
</div>
@@ -0,0 +1,18 @@
<script lang="ts">
interface Props {
currentItem: { name?: string } | null;
displayPreview: string | undefined;
}
let { currentItem, displayPreview }: Props = $props();
</script>
{#if displayPreview}
<div class="flex flex-1 items-center justify-center">
<img
src={displayPreview}
alt={currentItem?.name || 'preview'}
class="max-h-[80vh] max-w-[80vw] rounded-lg object-contain shadow-lg"
/>
</div>
{/if}
@@ -0,0 +1,174 @@
<script lang="ts">
import type { ChatAttachmentDisplayItem } from '$lib/types';
import { FileText, Eye, Info } from '@lucide/svelte';
import { Button } from '$lib/components/ui/button';
import * as Alert from '$lib/components/ui/alert';
import { SyntaxHighlightedCode } from '$lib/components/app';
import { getLanguageFromFilename } from '$lib/utils';
import { convertPDFToImage } from '$lib/utils/browser-only';
import { PdfViewMode } from '$lib/enums';
interface Props {
currentItem: ChatAttachmentDisplayItem | null;
displayName: string;
displayTextContent: string | undefined;
hasVisionModality: boolean;
activeModelId?: string;
}
let { currentItem, displayName, displayTextContent, hasVisionModality, activeModelId }: Props =
$props();
let pdfViewMode = $state<PdfViewMode>(PdfViewMode.PAGES);
let pdfImages = $state<string[]>([]);
let pdfImagesLoading = $state(false);
let pdfImagesError = $state<string | null>(null);
let language = $derived(getLanguageFromFilename(displayName));
async function loadPdfImages() {
if (pdfImages.length > 0 || pdfImagesLoading || !currentItem) return;
pdfImagesLoading = true;
pdfImagesError = null;
try {
let file: File | null = null;
if (currentItem.uploadedFile?.file) {
file = currentItem.uploadedFile.file;
} else if (currentItem.attachment) {
// Check if we have pre-processed images
if (
'images' in currentItem.attachment &&
currentItem.attachment.images &&
Array.isArray(currentItem.attachment.images) &&
currentItem.attachment.images.length > 0
) {
pdfImages = currentItem.attachment.images;
return;
}
// Convert base64 back to File for processing
if ('base64Data' in currentItem.attachment && currentItem.attachment.base64Data) {
const base64Data = currentItem.attachment.base64Data;
const byteCharacters = atob(base64Data);
const byteNumbers = new Array(byteCharacters.length);
for (let i = 0; i < byteCharacters.length; i++) {
byteNumbers[i] = byteCharacters.charCodeAt(i);
}
const byteArray = new Uint8Array(byteNumbers);
file = new File([byteArray], displayName, { type: 'application/pdf' });
}
}
if (file) {
pdfImages = await convertPDFToImage(file);
} else {
throw new Error('No PDF file available for conversion');
}
} catch (error) {
pdfImagesError = error instanceof Error ? error.message : 'Failed to load PDF images';
} finally {
pdfImagesLoading = false;
}
}
$effect(() => {
if (pdfViewMode === PdfViewMode.PAGES) {
loadPdfImages();
}
});
</script>
<div class="mb-4 flex items-center justify-end gap-2">
<Button
variant={pdfViewMode === PdfViewMode.TEXT ? 'default' : 'outline'}
size="sm"
onclick={() => (pdfViewMode = PdfViewMode.TEXT)}
disabled={pdfImagesLoading}
>
<FileText class="mr-1 h-4 w-4" />
Text
</Button>
<Button
variant={pdfViewMode === PdfViewMode.PAGES ? 'default' : 'outline'}
size="sm"
onclick={() => (pdfViewMode = PdfViewMode.PAGES)}
disabled={pdfImagesLoading}
>
{#if pdfImagesLoading}
<div
class="mr-1 h-4 w-4 animate-spin rounded-full border-2 border-current border-t-transparent"
></div>
{:else}
<Eye class="mr-1 h-4 w-4" />
{/if}
Pages
</Button>
</div>
{#if !hasVisionModality && activeModelId && currentItem}
<Alert.Root class="mb-4 max-w-4xl">
<Info class="h-4 w-4" />
<Alert.Title>Preview only</Alert.Title>
<Alert.Description>
<span class="inline-flex">
The selected model does not support vision. Only the extracted
<!-- svelte-ignore a11y_click_events_have_key_events -->
<!-- svelte-ignore a11y_no_static_element_interactions -->
<span
class="mx-1 cursor-pointer underline"
onclick={() => (pdfViewMode = PdfViewMode.TEXT)}
>
text
</span>
will be sent to the model.
</span>
</Alert.Description>
</Alert.Root>
{/if}
{#if pdfImagesLoading}
<div class="flex flex-1 items-center justify-center p-8">
<div class="text-center">
<div
class="mx-auto mb-4 h-8 w-8 animate-spin rounded-full border-4 border-white border-t-transparent"
></div>
<p class="text-white/70">Converting PDF to images...</p>
</div>
</div>
{:else if pdfImagesError}
<div class="flex flex-1 items-center justify-center p-8">
<div class="text-center">
<FileText class="mx-auto mb-4 h-16 w-16 text-white/50" />
<p class="mb-4 text-white/70">Failed to load PDF images</p>
<p class="text-sm text-white/50">{pdfImagesError}</p>
</div>
</div>
{:else if pdfImages.length > 0}
{#each pdfImages as image, index (image)}
<p class="mb-2 text-sm text-white/50">Page {index + 1}</p>
<img src={image} alt="PDF Page {index + 1}" class="mx-auto max-w-[85vw] rounded-lg shadow-lg" />
<div class="h-4"></div>
{/each}
{:else}
<div class="flex flex-1 items-center justify-center p-8">
<div class="text-center">
<FileText class="mx-auto mb-4 h-16 w-16 text-white/50" />
<p class="text-white/70">No PDF pages available</p>
</div>
</div>
{/if}
{#if pdfViewMode === PdfViewMode.TEXT && displayTextContent}
<div class="px-4 pb-4">
<SyntaxHighlightedCode
class="max-w-4xl"
code={displayTextContent}
{language}
maxHeight="none"
/>
</div>
{/if}
@@ -0,0 +1,21 @@
<script lang="ts">
import { SyntaxHighlightedCode } from '$lib/components/app';
interface Props {
displayTextContent: string | undefined;
language: string;
}
let { displayTextContent, language }: Props = $props();
</script>
{#if displayTextContent}
<div class="px-4 pb-4">
<SyntaxHighlightedCode
class="max-w-4xl"
code={displayTextContent}
{language}
maxHeight="none"
/>
</div>
{/if}
@@ -0,0 +1,17 @@
<script lang="ts">
import type { Component } from 'svelte';
interface Props {
IconComponent: Component;
}
let { IconComponent }: Props = $props();
</script>
<div class="flex flex-1 items-center justify-center p-8">
<div class="text-center">
<IconComponent class="mx-auto mb-4 h-16 w-16 text-white/50" />
<p class="text-white/70">Preview not available for this file type</p>
</div>
</div>
@@ -0,0 +1,16 @@
<script lang="ts">
interface Props {
displayName: string;
fileSize: string;
}
let { displayName, fileSize }: Props = $props();
</script>
<div class="sticky top-0 z-[20] mb-4 rounded-lg bg-black/5 px-4 py-2 text-center backdrop-blur-md">
<p class="font-medium text-white">{displayName}</p>
{#if fileSize}
<p class="text-xs text-white/60">{fileSize}</p>
{/if}
</div>
@@ -0,0 +1,34 @@
<script lang="ts">
import { ChevronLeft, ChevronRight } from '@lucide/svelte';
import { Button } from '$lib/components/ui/button';
interface Props {
onPrev: () => void;
onNext: () => void;
show: boolean;
}
let { onPrev, onNext, show }: Props = $props();
</script>
{#if show}
<Button
variant="secondary"
size="icon"
class="absolute top-1/2 left-4 z-10 h-8 w-8 -translate-y-1/2 rounded-full bg-background/5 p-0 text-white!"
onclick={onPrev}
aria-label="Previous"
>
<ChevronLeft class="size-4" />
</Button>
<Button
variant="secondary"
size="icon"
class="absolute top-1/2 right-4 z-10 h-8 w-8 -translate-y-1/2 rounded-full bg-background/5 p-0 text-white!"
onclick={onNext}
aria-label="Next"
>
<ChevronRight class="size-4" />
</Button>
{/if}
@@ -0,0 +1,63 @@
<script lang="ts">
import { Music, FileText } from '@lucide/svelte';
import { HorizontalScrollCarousel } from '$lib/components/app/misc';
interface PreviewItem {
id: string;
name: string;
isImage: boolean;
isAudio: boolean;
preview?: string;
}
interface Props {
items: PreviewItem[];
currentIndex: number;
onNavigate: (index: number) => void;
}
let { items, currentIndex, onNavigate }: Props = $props();
function getFileExtension(name: string): string {
const parts = name.split('.');
if (parts.length > 1) {
return parts.pop()?.toUpperCase() ?? '';
}
return '';
}
</script>
{#if items.length > 1}
<div class="sticky bottom-0 z-10 mt-4 flex-shrink-0">
<HorizontalScrollCarousel class="max-w-full">
{#each items as item, index (item.id)}
<button
data-thumbnail-index={index}
class={[
'relative flex-shrink-0 cursor-pointer overflow-hidden rounded border-2 bg-black/80 backdrop-blur-sm transition-all hover:opacity-90',
index === currentIndex ? 'border-white' : 'border-transparent opacity-60',
'[&:not(:first-child)]:last:mr-4 [&:not(:last-child)]:first:ml-4'
]}
onclick={() => onNavigate(index)}
aria-label={`Go to ${item.name}`}
>
{#if item.isImage && item.preview}
<img src={item.preview} alt={item.name} class="h-12 w-12 object-cover" />
{:else}
<div
class="bg-foreground-muted/50 flex h-12 w-12 flex-col items-center justify-center gap-0.5 py-1"
>
{#if item.isAudio}
<Music class="h-4 w-4 text-white/70" />
{:else}
<FileText class="h-4 w-4 text-white/70" />
{/if}
<span class="font-mono text-[9px] text-white/60">{getFileExtension(item.name)}</span>
</div>
{/if}
</button>
{/each}
</HorizontalScrollCarousel>
</div>
{/if}
@@ -1,117 +0,0 @@
<script lang="ts">
import {
ChatAttachmentThumbnailImage,
ChatAttachmentThumbnailFile,
DialogChatAttachmentPreview
} from '$lib/components/app';
import { getAttachmentDisplayItems } from '$lib/utils';
interface Props {
uploadedFiles?: ChatUploadedFile[];
attachments?: DatabaseMessageExtra[];
readonly?: boolean;
onFileRemove?: (fileId: string) => void;
imageHeight?: string;
imageWidth?: string;
imageClass?: string;
activeModelId?: string;
}
let {
uploadedFiles = [],
attachments = [],
readonly = false,
onFileRemove,
imageHeight = 'h-24',
imageWidth = 'w-auto',
imageClass = '',
activeModelId
}: Props = $props();
let previewDialogOpen = $state(false);
let previewItem = $state<ChatAttachmentPreviewItem | null>(null);
let displayItems = $derived(getAttachmentDisplayItems({ uploadedFiles, attachments }));
let imageItems = $derived(displayItems.filter((item) => item.isImage));
let fileItems = $derived(displayItems.filter((item) => !item.isImage));
function openPreview(item: (typeof displayItems)[0], event?: Event) {
if (event) {
event.preventDefault();
event.stopPropagation();
}
previewItem = {
uploadedFile: item.uploadedFile,
attachment: item.attachment,
preview: item.preview,
name: item.name,
size: item.size,
textContent: item.textContent
};
previewDialogOpen = true;
}
</script>
<div class="space-y-4">
<div class="min-h-0 flex-1 space-y-6 overflow-y-auto px-1">
{#if fileItems.length > 0}
<div>
<h3 class="mb-3 text-sm font-medium text-foreground">Files ({fileItems.length})</h3>
<div class="flex flex-wrap items-start gap-3">
{#each fileItems as item (item.id)}
<ChatAttachmentThumbnailFile
class="cursor-pointer"
id={item.id}
name={item.name}
size={item.size}
{readonly}
onRemove={onFileRemove}
textContent={item.textContent}
attachment={item.attachment}
uploadedFile={item.uploadedFile}
onClick={(event?: MouseEvent) => openPreview(item, event)}
/>
{/each}
</div>
</div>
{/if}
{#if imageItems.length > 0}
<div>
<h3 class="mb-3 text-sm font-medium text-foreground">Images ({imageItems.length})</h3>
<div class="flex flex-wrap items-start gap-3">
{#each imageItems as item (item.id)}
{#if item.preview}
<ChatAttachmentThumbnailImage
class="cursor-pointer"
id={item.id}
name={item.name}
preview={item.preview}
{readonly}
onRemove={onFileRemove}
height={imageHeight}
width={imageWidth}
{imageClass}
onClick={(event) => openPreview(item, event)}
/>
{/if}
{/each}
</div>
</div>
{/if}
</div>
</div>
{#if previewItem}
<DialogChatAttachmentPreview
bind:open={previewDialogOpen}
uploadedFile={previewItem.uploadedFile}
attachment={previewItem.attachment}
preview={previewItem.preview}
name={previewItem.name}
size={previewItem.size}
textContent={previewItem.textContent}
{activeModelId}
/>
{/if}
@@ -1,14 +1,13 @@
<script lang="ts">
import {
ChatAttachmentsList,
ChatAttachmentMcpResources,
ChatFormActions,
ChatFormFileInputInvisible,
ChatFormPromptPicker,
ChatFormResourcePicker,
ChatFormTextarea
ChatFormMcpResourcesList,
ChatFormPickers,
ChatFormTextarea,
DialogMcpResourcesBrowser
} from '$lib/components/app';
import { DialogMcpResources } from '$lib/components/app/dialogs';
import {
CLIPBOARD_CONTENT_QUOTE_PREFIX,
INPUT_CLASSES,
@@ -54,6 +53,8 @@
isLoading?: boolean;
placeholder?: string;
showMcpPromptButton?: boolean;
showAddButton?: boolean;
showModelSelector?: boolean;
// Event Handlers
onAttachmentRemove?: (index: number) => void;
@@ -73,6 +74,8 @@
isLoading = false,
placeholder = 'Type a message...',
showMcpPromptButton = false,
showAddButton = true,
showModelSelector = true,
uploadedFiles = $bindable([]),
value = $bindable(''),
onAttachmentRemove,
@@ -85,31 +88,21 @@
onValueChange
}: Props = $props();
/**
*
*
* STATE
*
*
*/
// Component References
let audioRecorder: AudioRecorder | undefined;
let chatFormActionsRef: ChatFormActions | undefined = $state(undefined);
let fileInputRef: ChatFormFileInputInvisible | undefined = $state(undefined);
let promptPickerRef: ChatFormPromptPicker | undefined = $state(undefined);
let resourcePickerRef: ChatFormResourcePicker | undefined = $state(undefined);
let pickersRef: { handleKeydown: (event: KeyboardEvent) => boolean } | undefined =
$state(undefined);
let textareaRef: ChatFormTextarea | undefined = $state(undefined);
// Audio Recording State
let isRecording = $state(false);
let recordingSupported = $state(false);
// Prompt Picker State
// Picker State
let isPromptPickerOpen = $state(false);
let promptSearchQuery = $state('');
// Inline Resource Picker State (triggered by @)
let isInlineResourcePickerOpen = $state(false);
let resourceSearchQuery = $state('');
@@ -117,22 +110,12 @@
let isResourceDialogOpen = $state(false);
let preSelectedResourceUri = $state<string | undefined>(undefined);
/**
*
*
* DERIVED STATE
*
*
*/
// Configuration
let currentConfig = $derived(config());
let pasteLongTextToFileLength = $derived.by(() => {
const n = Number(currentConfig.pasteLongTextToFileLen);
return Number.isNaN(n) ? Number(SETTING_CONFIG_DEFAULT.pasteLongTextToFileLen) : n;
});
// Model Selection Logic
let isRouter = $derived(isRouterMode());
let conversationModel = $derived(
chatStore.getConversationModel(activeMessages() as DatabaseMessage[])
@@ -158,7 +141,6 @@
return null;
});
// Form Validation State
let hasModelSelected = $derived(!isRouter || !!conversationModel || !!selectedModelId());
let hasLoadingAttachments = $derived(uploadedFiles.some((f) => f.isLoading));
let hasAttachments = $derived(
@@ -166,27 +148,11 @@
);
let canSubmit = $derived(value.trim().length > 0 || hasAttachments);
/**
*
*
* LIFECYCLE
*
*
*/
onMount(() => {
recordingSupported = isAudioRecordingSupported();
audioRecorder = new AudioRecorder();
});
/**
*
*
* PUBLIC API
*
*
*/
export function focus() {
textareaRef?.focus();
}
@@ -199,10 +165,6 @@
chatFormActionsRef?.openModelSelector();
}
/**
* Check if a model is selected, open selector if not
* @returns true if model is selected, false otherwise
*/
export function checkModelSelected(): boolean {
if (!hasModelSelected) {
chatFormActionsRef?.openModelSelector();
@@ -211,14 +173,6 @@
return true;
}
/**
*
*
* EVENT HANDLERS - File Management
*
*
*/
function handleFileSelect(files: File[]) {
onFilesAdd?.(files);
}
@@ -238,14 +192,6 @@
}
}
/**
*
*
* EVENT HANDLERS - Input & Keyboard
*
*
*/
function handleInput() {
const perChatOverrides = conversationsStore.getAllMcpServerOverrides();
const hasServers = mcpStore.hasEnabledServers(perChatOverrides);
@@ -273,11 +219,7 @@
}
function handleKeydown(event: KeyboardEvent) {
if (isPromptPickerOpen && promptPickerRef?.handleKeydown(event)) {
return;
}
if (isInlineResourcePickerOpen && resourcePickerRef?.handleKeydown(event)) {
if (pickersRef?.handleKeydown(event)) {
return;
}
@@ -388,14 +330,6 @@
}
}
/**
*
*
* EVENT HANDLERS - Prompt Picker
*
*
*/
function handlePromptLoadStart(
placeholderId: string,
promptInfo: MCPPromptInfo,
@@ -474,14 +408,6 @@
textareaRef?.focus();
}
/**
*
*
* EVENT HANDLERS - Inline Resource Picker
*
*
*/
function handleInlineResourcePickerClose() {
isInlineResourcePickerOpen = false;
resourceSearchQuery = '';
@@ -489,7 +415,6 @@
}
function handleInlineResourceSelect() {
// Clear the @query from input after resource is attached
if (value.startsWith(RESOURCE_TRIGGER_PREFIX)) {
value = '';
onValueChange?.('');
@@ -512,14 +437,6 @@
isResourceDialogOpen = true;
}
/**
*
*
* EVENT HANDLERS - Audio Recording
*
*
*/
async function handleMicClick() {
if (!audioRecorder || !recordingSupported) {
console.warn('Audio recording not supported');
@@ -552,29 +469,27 @@
<form
class="relative {className}"
onsubmit={(e) => {
e.preventDefault();
onsubmit={(event) => {
event.preventDefault();
if (!canSubmit || disabled || hasLoadingAttachments) return;
onSubmit?.();
}}
>
<ChatFormPromptPicker
bind:this={promptPickerRef}
isOpen={isPromptPickerOpen}
searchQuery={promptSearchQuery}
onClose={handlePromptPickerClose}
<ChatFormPickers
bind:this={pickersRef}
{isPromptPickerOpen}
{promptSearchQuery}
{isInlineResourcePickerOpen}
{resourceSearchQuery}
onPromptPickerClose={handlePromptPickerClose}
onInlineResourcePickerClose={handleInlineResourcePickerClose}
onInlineResourceSelect={handleInlineResourceSelect}
onPromptLoadStart={handlePromptLoadStart}
onPromptLoadComplete={handlePromptLoadComplete}
onPromptLoadError={handlePromptLoadError}
/>
<ChatFormResourcePicker
bind:this={resourcePickerRef}
isOpen={isInlineResourcePickerOpen}
searchQuery={resourceSearchQuery}
onClose={handleInlineResourcePickerClose}
onResourceSelect={handleInlineResourceSelect}
onBrowse={handleBrowseResources}
onInlineResourceBrowse={handleBrowseResources}
/>
<div
@@ -611,7 +526,7 @@
/>
{#if mcpHasResourceAttachments()}
<ChatAttachmentMcpResources
<ChatFormMcpResourcesList
class="mb-3"
onResourceClick={(uri) => {
preSelectedResourceUri = uri;
@@ -624,10 +539,11 @@
class="px-3"
bind:this={chatFormActionsRef}
canSend={canSubmit}
hasText={value.trim().length > 0}
{disabled}
{isLoading}
{isRecording}
{showAddButton}
{showModelSelector}
{uploadedFiles}
onFileUpload={handleFileUpload}
onMicClick={handleMicClick}
@@ -640,7 +556,7 @@
</div>
</form>
<DialogMcpResources
<DialogMcpResourcesBrowser
bind:open={isResourceDialogOpen}
preSelectedUri={preSelectedResourceUri}
onAttach={(resource: MCPResourceInfo) => {
@@ -0,0 +1,33 @@
<script lang="ts">
import { Plus } from '@lucide/svelte';
import { Button } from '$lib/components/ui/button';
import * as Tooltip from '$lib/components/ui/tooltip';
import { ATTACHMENT_TOOLTIP_TEXT } from '$lib/constants';
interface Props {
disabled?: boolean;
onclick?: (e: MouseEvent) => void;
}
let { disabled = false, onclick }: Props = $props();
</script>
<Tooltip.Root>
<Tooltip.Trigger class="w-full">
<Button
class="file-upload-button h-8 w-8 rounded-full p-0"
{disabled}
{onclick}
variant="secondary"
type="button"
>
<span class="sr-only">{ATTACHMENT_TOOLTIP_TEXT}</span>
<Plus class="h-4 w-4" />
</Button>
</Tooltip.Trigger>
<Tooltip.Content>
<p>{ATTACHMENT_TOOLTIP_TEXT}</p>
</Tooltip.Content>
</Tooltip.Root>
@@ -1,17 +1,18 @@
<script lang="ts">
import { Plus } from '@lucide/svelte';
import { Button } from '$lib/components/ui/button';
import type { Snippet } from 'svelte';
import * as DropdownMenu from '$lib/components/ui/dropdown-menu';
import * as Tooltip from '$lib/components/ui/tooltip';
import {
ATTACHMENT_FILE_ITEMS,
ATTACHMENT_EXTRA_ITEMS,
ATTACHMENT_MCP_ITEMS,
ATTACHMENT_TOOLTIP_TEXT,
TOOLTIP_DELAY_DURATION
} from '$lib/constants';
import { AttachmentMenuItemId } from '$lib/enums';
import { ChatFormActionToolsSubmenu, ChatFormActionMcpServersSubmenu } from '$lib/components/app';
import {
ChatFormActionAddToolsSubmenu,
ChatFormActionAddMcpServersSubmenu
} from '$lib/components/app';
import { useAttachmentMenu } from '$lib/hooks/use-attachment-menu.svelte';
@@ -27,6 +28,7 @@
onMcpPromptClick?: () => void;
onMcpSettingsClick?: () => void;
onMcpResourcesClick?: () => void;
trigger: Snippet<[{ disabled: boolean }]>;
}
let {
@@ -40,7 +42,8 @@
onSystemPromptClick,
onMcpPromptClick,
onMcpSettingsClick,
onMcpResourcesClick
onMcpResourcesClick,
trigger
}: Props = $props();
let dropdownOpen = $state(false);
@@ -62,24 +65,7 @@
<div class="flex items-center gap-1 {className}">
<DropdownMenu.Root bind:open={dropdownOpen}>
<DropdownMenu.Trigger name="Attach files" {disabled}>
<Tooltip.Root>
<Tooltip.Trigger class="w-full">
<Button
class="file-upload-button h-8 w-8 rounded-full p-0"
{disabled}
variant="secondary"
type="button"
>
<span class="sr-only">{ATTACHMENT_TOOLTIP_TEXT}</span>
<Plus class="h-4 w-4" />
</Button>
</Tooltip.Trigger>
<Tooltip.Content>
<p>{ATTACHMENT_TOOLTIP_TEXT}</p>
</Tooltip.Content>
</Tooltip.Root>
{@render trigger({ disabled })}
</DropdownMenu.Trigger>
<DropdownMenu.Content align="start" class="w-48">
@@ -161,9 +147,9 @@
{/if}
{/each}
<ChatFormActionToolsSubmenu />
<ChatFormActionAddToolsSubmenu />
<ChatFormActionMcpServersSubmenu onMcpSettingsClick={handleMcpSettingsClick} />
<ChatFormActionAddMcpServersSubmenu onMcpSettingsClick={handleMcpSettingsClick} />
{#each ATTACHMENT_MCP_ITEMS as item (item.id)}
{#if attachmentMenu.isItemVisible(item.visibleWhen)}
@@ -0,0 +1,149 @@
<script lang="ts">
import { Settings, Plus } from '@lucide/svelte';
import { Switch } from '$lib/components/ui/switch';
import * as DropdownMenu from '$lib/components/ui/dropdown-menu';
import { McpLogo, DropdownMenuSearchable } from '$lib/components/app';
import { conversationsStore } from '$lib/stores/conversations.svelte';
import { mcpStore } from '$lib/stores/mcp.svelte';
import { HealthCheckStatus } from '$lib/enums';
import type { MCPServerSettingsEntry } from '$lib/types';
import { goto } from '$app/navigation';
interface Props {
onMcpSettingsClick?: () => void;
}
let { onMcpSettingsClick }: Props = $props();
let mcpSearchQuery = $state('');
let allMcpServers = $derived(mcpStore.getServersSorted());
let mcpServers = $derived(allMcpServers.filter((s) => s.enabled));
let hasMcpServers = $derived(mcpServers.length > 0);
// let hasAnyMcpServers = $derived(allMcpServers.length > 0);
let filteredMcpServers = $derived.by(() => {
const query = mcpSearchQuery.toLowerCase().trim();
if (!query) return mcpServers;
return mcpServers.filter((s) => {
const name = getServerLabel(s).toLowerCase();
const url = s.url.toLowerCase();
return name.includes(query) || url.includes(query);
});
});
function getServerLabel(server: MCPServerSettingsEntry): string {
return mcpStore.getServerLabel(server);
}
function isServerEnabledForChat(serverId: string): boolean {
return conversationsStore.isMcpServerEnabledForChat(serverId);
}
async function toggleServerForChat(serverId: string) {
await conversationsStore.toggleMcpServerForChat(serverId);
}
function handleMcpSubMenuOpen(open: boolean) {
if (open) {
mcpSearchQuery = '';
mcpStore.runHealthChecksForServers(allMcpServers);
}
}
function handleMcpSettingsClick() {
onMcpSettingsClick?.();
goto(`${hasMcpServers ? '' : '?add'}#/settings/mcp`);
}
</script>
<DropdownMenu.Root>
<DropdownMenu.Sub onOpenChange={handleMcpSubMenuOpen}>
<DropdownMenu.SubTrigger class="flex cursor-pointer items-center gap-2">
<McpLogo class="h-4 w-4" />
<span>MCP Servers</span>
</DropdownMenu.SubTrigger>
<DropdownMenu.SubContent class="w-72 pt-0">
{#if hasMcpServers}
<DropdownMenuSearchable
placeholder="Search servers..."
bind:searchValue={mcpSearchQuery}
emptyMessage="No servers found"
isEmpty={filteredMcpServers.length === 0}
>
<div class="max-h-64 overflow-y-auto">
{#each filteredMcpServers as server (server.id)}
{@const healthState = mcpStore.getHealthCheckState(server.id)}
{@const hasError = healthState.status === HealthCheckStatus.ERROR}
{@const isEnabledForChat = isServerEnabledForChat(server.id)}
<button
type="button"
class="flex w-full items-center justify-between gap-2 rounded-sm px-2 py-2 text-left transition-colors hover:bg-accent disabled:cursor-not-allowed disabled:opacity-50"
onclick={() => !hasError && toggleServerForChat(server.id)}
disabled={hasError}
>
<div class="flex min-w-0 flex-1 items-center gap-2">
{#if mcpStore.getServerFavicon(server.id)}
<img
src={mcpStore.getServerFavicon(server.id)}
alt=""
class="h-4 w-4 shrink-0 rounded-sm"
onerror={(e) => {
(e.currentTarget as HTMLImageElement).style.display = 'none';
}}
/>
{/if}
<span class="truncate text-sm">{getServerLabel(server)}</span>
{#if hasError}
<span
class="shrink-0 rounded bg-destructive/15 px-1.5 py-0.5 text-xs text-destructive"
>
Error
</span>
{/if}
</div>
<Switch
checked={isEnabledForChat}
disabled={hasError}
onclick={(e) => e.stopPropagation()}
onCheckedChange={() => toggleServerForChat(server.id)}
/>
</button>
{/each}
</div>
{#snippet footer()}
<DropdownMenu.Item
class="flex cursor-pointer items-center gap-2"
onclick={handleMcpSettingsClick}
>
<Settings class="h-4 w-4" />
<span>Manage MCP Servers</span>
</DropdownMenu.Item>
{/snippet}
</DropdownMenuSearchable>
{:else}
<div class="px-2 py-3 text-center text-sm text-muted-foreground">
No MCP servers configured
</div>
<DropdownMenu.Separator />
<DropdownMenu.Item
class="flex cursor-pointer items-center gap-2"
onclick={handleMcpSettingsClick}
>
<Plus class="h-4 w-4" />
<span>Add MCP Servers</span>
</DropdownMenu.Item>
{/if}
</DropdownMenu.SubContent>
</DropdownMenu.Sub>
</DropdownMenu.Root>
@@ -1,18 +1,17 @@
<script lang="ts">
import { Plus } from '@lucide/svelte';
import { Button } from '$lib/components/ui/button';
import type { Snippet } from 'svelte';
import * as Tooltip from '$lib/components/ui/tooltip';
import * as Sheet from '$lib/components/ui/sheet';
import { TOOLTIP_DELAY_DURATION } from '$lib/constants';
import {
ATTACHMENT_FILE_ITEMS,
ATTACHMENT_EXTRA_ITEMS,
ATTACHMENT_MCP_ITEMS,
ATTACHMENT_TOOLTIP_TEXT
ATTACHMENT_MCP_ITEMS
} from '$lib/constants/attachment-menu';
import { ChatFormActionToolsSubmenu, ChatFormActionMcpServersSubmenu } from '$lib/components/app';
import { McpLogo } from '$lib/components/app';
import { useAttachmentMenu } from '$lib/hooks/use-attachment-menu.svelte';
import { AttachmentMenuItemId } from '$lib/enums';
import { PencilRuler } from '@lucide/svelte';
interface Props {
class?: string;
@@ -24,8 +23,8 @@
onFileUpload?: () => void;
onSystemPromptClick?: () => void;
onMcpPromptClick?: () => void;
onMcpSettingsClick?: () => void;
onMcpResourcesClick?: () => void;
trigger: Snippet<[{ disabled: boolean; onclick?: () => void }]>;
}
let {
@@ -38,8 +37,8 @@
onFileUpload,
onSystemPromptClick,
onMcpPromptClick,
onMcpSettingsClick,
onMcpResourcesClick
onMcpResourcesClick,
trigger
}: Props = $props();
let sheetOpen = $state(false);
@@ -52,28 +51,14 @@
}
);
function handleMcpSettingsClick() {
sheetOpen = false;
onMcpSettingsClick?.();
}
const sheetItemClass =
'flex w-full items-center gap-3 rounded-md px-3 py-2.5 text-left text-sm transition-colors hover:bg-accent active:bg-accent disabled:cursor-not-allowed disabled:opacity-50';
</script>
<div class="flex items-center gap-1 {className}">
<Sheet.Root bind:open={sheetOpen}>
<Button
class="file-upload-button h-8 w-8 rounded-full p-0"
{disabled}
variant="secondary"
type="button"
onclick={() => (sheetOpen = true)}
>
<span class="sr-only">{ATTACHMENT_TOOLTIP_TEXT}</span>
<Plus class="h-4 w-4" />
</Button>
{@render trigger({ disabled, onclick: () => (sheetOpen = true) })}
<!-- <ChatFormActionAddButton {disabled} onclick={() => (sheetOpen = true)} /> -->
<Sheet.Content side="bottom" class="max-h-[85vh] gap-0 overflow-y-auto">
<Sheet.Header>
@@ -161,9 +146,17 @@
<div class="my-2 border-t"></div>
<ChatFormActionToolsSubmenu />
<a href="#/settings/mcp" class="flex items-center gap-3 px-3 py-2">
<McpLogo class="inline h-4 w-4" />
<ChatFormActionMcpServersSubmenu onMcpSettingsClick={handleMcpSettingsClick} />
<span class="text-sm">MCP Servers</span>
</a>
<a href="#/settings/chat/tools" class="flex items-center gap-3 px-3 py-2">
<PencilRuler class="inline h-4 w-4" />
<span class="text-sm">Tools</span>
</a>
{#each ATTACHMENT_MCP_ITEMS as item (item.id)}
{#if attachmentMenu.isItemVisible(item.visibleWhen)}
@@ -24,6 +24,7 @@
{#if toolsStore.loading}
<div class="px-3 py-4 text-center text-sm text-muted-foreground">
<Loader2 class="mx-auto mb-1 h-4 w-4 animate-spin" />
Loading tools...
</div>
{:else if toolsStore.isToolsEndpointUnreachable}
@@ -31,19 +32,21 @@
<span class="flex gap-2">
<Info class="mt-0.5 h-4 w-4 shrink-0" />
<span
>Run llama-server with <code>--tools</code> flag to enable
<strong>Built-in Tools</strong>.</span
>
<span>
Run llama-server with <code>--tools</code> flag to enable
<strong>Built-in Tools</strong>.
</span>
</span>
<span class="flex gap-2">
<Info class="mt-0.5 h-4 w-4 shrink-0" />
<span
>{hasMcpServersAvailable ? 'Enable' : 'Add'} MCP Server(s) to access
<strong>MCP Tools</strong>.</span
>
<span>
{hasMcpServersAvailable ? 'Enable' : 'Add'} MCP Server(s) to access
<strong>MCP Tools</strong>.
</span>
</span>
</div>
{:else if toolsStore.error}
@@ -0,0 +1,68 @@
<script lang="ts">
import { IsMobile } from '$lib/hooks/is-mobile.svelte';
import ChatFormActionAddDropdown from './ChatFormActionAddDropdown.svelte';
import ChatFormActionAddSheet from './ChatFormActionAddSheet.svelte';
import ChatFormActionAddButton from './ChatFormActionAddButton.svelte';
interface Props {
disabled?: boolean;
hasAudioModality?: boolean;
hasMcpPromptsSupport?: boolean;
hasMcpResourcesSupport?: boolean;
hasVisionModality?: boolean;
onFileUpload?: () => void;
onMcpPromptClick?: () => void;
onMcpResourcesClick?: () => void;
onMcpSettingsClick?: () => void;
onSystemPromptClick?: () => void;
}
let {
disabled = false,
hasAudioModality = false,
hasMcpPromptsSupport = false,
hasMcpResourcesSupport = false,
hasVisionModality = false,
onFileUpload,
onMcpPromptClick,
onMcpResourcesClick,
onMcpSettingsClick,
onSystemPromptClick
}: Props = $props();
const isMobile = new IsMobile();
</script>
{#if isMobile.current}
<ChatFormActionAddSheet
{disabled}
{hasAudioModality}
{hasVisionModality}
{hasMcpPromptsSupport}
{hasMcpResourcesSupport}
{onFileUpload}
{onMcpPromptClick}
{onMcpResourcesClick}
>
{#snippet trigger({ disabled, onclick })}
<ChatFormActionAddButton {disabled} {onclick} />
{/snippet}
</ChatFormActionAddSheet>
{:else}
<ChatFormActionAddDropdown
{disabled}
{hasAudioModality}
{hasVisionModality}
{hasMcpPromptsSupport}
{hasMcpResourcesSupport}
{onFileUpload}
{onMcpPromptClick}
{onMcpResourcesClick}
{onMcpSettingsClick}
{onSystemPromptClick}
>
{#snippet trigger()}
<ChatFormActionAddButton {disabled} />
{/snippet}
</ChatFormActionAddDropdown>
{/if}
@@ -1,147 +0,0 @@
<script lang="ts">
import { Settings, Plus } from '@lucide/svelte';
import { Switch } from '$lib/components/ui/switch';
import * as DropdownMenu from '$lib/components/ui/dropdown-menu';
import { McpLogo, DropdownMenuSearchable } from '$lib/components/app';
import { conversationsStore } from '$lib/stores/conversations.svelte';
import { mcpStore } from '$lib/stores/mcp.svelte';
import { HealthCheckStatus } from '$lib/enums';
import type { MCPServerSettingsEntry } from '$lib/types';
import { goto } from '$app/navigation';
interface Props {
onMcpSettingsClick?: () => void;
}
let { onMcpSettingsClick }: Props = $props();
let mcpSearchQuery = $state('');
let allMcpServers = $derived(mcpStore.getServersSorted());
let mcpServers = $derived(allMcpServers.filter((s) => s.enabled));
let hasMcpServers = $derived(mcpServers.length > 0);
// let hasAnyMcpServers = $derived(allMcpServers.length > 0);
let filteredMcpServers = $derived.by(() => {
const query = mcpSearchQuery.toLowerCase().trim();
if (!query) return mcpServers;
return mcpServers.filter((s) => {
const name = getServerLabel(s).toLowerCase();
const url = s.url.toLowerCase();
return name.includes(query) || url.includes(query);
});
});
function getServerLabel(server: MCPServerSettingsEntry): string {
return mcpStore.getServerLabel(server);
}
function isServerEnabledForChat(serverId: string): boolean {
return conversationsStore.isMcpServerEnabledForChat(serverId);
}
async function toggleServerForChat(serverId: string) {
await conversationsStore.toggleMcpServerForChat(serverId);
}
function handleMcpSubMenuOpen(open: boolean) {
if (open) {
mcpSearchQuery = '';
mcpStore.runHealthChecksForServers(allMcpServers);
}
}
function handleMcpSettingsClick() {
onMcpSettingsClick?.();
goto(`${hasMcpServers ? '' : '?add'}#/settings/mcp`);
}
</script>
<DropdownMenu.Sub onOpenChange={handleMcpSubMenuOpen}>
<DropdownMenu.SubTrigger class="flex cursor-pointer items-center gap-2">
<McpLogo class="h-4 w-4" />
<span>MCP Servers</span>
</DropdownMenu.SubTrigger>
<DropdownMenu.SubContent class="w-72 pt-0">
{#if hasMcpServers}
<DropdownMenuSearchable
placeholder="Search servers..."
bind:searchValue={mcpSearchQuery}
emptyMessage="No servers found"
isEmpty={filteredMcpServers.length === 0}
>
<div class="max-h-64 overflow-y-auto">
{#each filteredMcpServers as server (server.id)}
{@const healthState = mcpStore.getHealthCheckState(server.id)}
{@const hasError = healthState.status === HealthCheckStatus.ERROR}
{@const isEnabledForChat = isServerEnabledForChat(server.id)}
<button
type="button"
class="flex w-full items-center justify-between gap-2 rounded-sm px-2 py-2 text-left transition-colors hover:bg-accent disabled:cursor-not-allowed disabled:opacity-50"
onclick={() => !hasError && toggleServerForChat(server.id)}
disabled={hasError}
>
<div class="flex min-w-0 flex-1 items-center gap-2">
{#if mcpStore.getServerFavicon(server.id)}
<img
src={mcpStore.getServerFavicon(server.id)}
alt=""
class="h-4 w-4 shrink-0 rounded-sm"
onerror={(e) => {
(e.currentTarget as HTMLImageElement).style.display = 'none';
}}
/>
{/if}
<span class="truncate text-sm">{getServerLabel(server)}</span>
{#if hasError}
<span
class="shrink-0 rounded bg-destructive/15 px-1.5 py-0.5 text-xs text-destructive"
>
Error
</span>
{/if}
</div>
<Switch
checked={isEnabledForChat}
disabled={hasError}
onclick={(e: MouseEvent) => e.stopPropagation()}
onCheckedChange={() => toggleServerForChat(server.id)}
/>
</button>
{/each}
</div>
{#snippet footer()}
<DropdownMenu.Item
class="flex cursor-pointer items-center gap-2"
onclick={handleMcpSettingsClick}
>
<Settings class="h-4 w-4" />
<span>Manage MCP Servers</span>
</DropdownMenu.Item>
{/snippet}
</DropdownMenuSearchable>
{:else}
<div class="px-2 py-3 text-center text-sm text-muted-foreground">
No MCP servers configured
</div>
<DropdownMenu.Separator />
<DropdownMenu.Item
class="flex cursor-pointer items-center gap-2"
onclick={handleMcpSettingsClick}
>
<Plus class="h-4 w-4" />
<span>Add MCP Servers</span>
</DropdownMenu.Item>
{/if}
</DropdownMenu.SubContent>
</DropdownMenu.Sub>
@@ -0,0 +1,160 @@
<script lang="ts">
import { chatStore } from '$lib/stores/chat.svelte';
import { modelsStore, modelOptions, selectedModelId } from '$lib/stores/models.svelte';
import { isRouterMode, serverError } from '$lib/stores/server.svelte';
import { ModelsSelectorDropdown, ModelsSelectorSheet } from '$lib/components/app';
import { IsMobile } from '$lib/hooks/is-mobile.svelte';
import { activeMessages } from '$lib/stores/conversations.svelte';
interface Props {
currentModel?: string;
disabled?: boolean;
forceForegroundText?: boolean;
hasAudioModality?: boolean;
hasVisionModality?: boolean;
hasModelSelected?: boolean;
isSelectedModelInCache?: boolean;
submitTooltip?: string;
useGlobalSelection?: boolean;
}
let {
currentModel,
disabled = false,
forceForegroundText = false,
hasAudioModality = $bindable(false),
hasVisionModality = $bindable(false),
hasModelSelected = $bindable(false),
isSelectedModelInCache = $bindable(true),
submitTooltip = $bindable(''),
useGlobalSelection = false
}: Props = $props();
let isRouter = $derived(isRouterMode());
let isOffline = $derived(!!serverError());
let conversationModel = $derived(
chatStore.getConversationModel(activeMessages() as DatabaseMessage[])
);
let lastSyncedConversationModel: string | null = null;
$effect(() => {
if (conversationModel && conversationModel !== lastSyncedConversationModel) {
lastSyncedConversationModel = conversationModel;
modelsStore.selectModelByName(conversationModel);
} else if (isRouter && !modelsStore.selectedModelId && modelsStore.loadedModelIds.length > 0) {
lastSyncedConversationModel = null;
// auto-select the first loaded model only when nothing is selected yet
const first = modelOptions().find((m) => modelsStore.loadedModelIds.includes(m.model));
if (first) modelsStore.selectModelById(first.id);
}
});
let activeModelId = $derived.by(() => {
const options = modelOptions();
if (!isRouter) {
return options.length > 0 ? options[0].model : null;
}
const selectedId = selectedModelId();
if (selectedId) {
const model = options.find((m) => m.id === selectedId);
if (model) return model.model;
}
if (conversationModel) {
const model = options.find((m) => m.model === conversationModel);
if (model) return model.model;
}
return null;
});
let modelPropsVersion = $state(0); // Used to trigger reactivity after fetch
$effect(() => {
if (activeModelId) {
const cached = modelsStore.getModelProps(activeModelId);
if (!cached) {
modelsStore.fetchModelProps(activeModelId).then(() => {
modelPropsVersion++;
});
}
}
});
$effect(() => {
hasAudioModality = activeModelId ? modelsStore.modelSupportsAudio(activeModelId) : false;
});
$effect(() => {
void modelPropsVersion;
hasVisionModality = activeModelId ? modelsStore.modelSupportsVision(activeModelId) : false;
});
$effect(() => {
hasModelSelected = !isRouter || !!conversationModel || !!selectedModelId();
});
$effect(() => {
if (!isRouter) {
isSelectedModelInCache = true;
} else if (conversationModel) {
isSelectedModelInCache = modelOptions().some((option) => option.model === conversationModel);
} else {
const currentModelId = selectedModelId();
if (!currentModelId) {
isSelectedModelInCache = false;
} else {
isSelectedModelInCache = modelOptions().some((option) => option.id === currentModelId);
}
}
});
$effect(() => {
if (!hasModelSelected) {
submitTooltip = 'Please select a model first';
} else if (!isSelectedModelInCache) {
submitTooltip = 'Selected model is not available, please select another';
} else {
submitTooltip = '';
}
});
let selectorModelRef: ModelsSelectorDropdown | ModelsSelectorSheet | undefined =
$state(undefined);
let isMobile = new IsMobile();
export function open() {
selectorModelRef?.open();
}
</script>
{#if isMobile.current}
<ModelsSelectorSheet
disabled={disabled || isOffline}
bind:this={selectorModelRef}
{currentModel}
{forceForegroundText}
{useGlobalSelection}
/>
{:else}
<ModelsSelectorDropdown
disabled={disabled || isOffline}
bind:this={selectorModelRef}
{currentModel}
{forceForegroundText}
{useGlobalSelection}
/>
{/if}
@@ -2,7 +2,6 @@
import { ArrowUp } from '@lucide/svelte';
import { Button } from '$lib/components/ui/button';
import * as Tooltip from '$lib/components/ui/tooltip';
import { cn } from '$lib/components/ui/utils';
interface Props {
canSend?: boolean;
@@ -20,12 +19,11 @@
<Button
type="submit"
disabled={isDisabled}
class={cn(
class={[
'h-8 w-8 rounded-full p-0',
showErrorState
? 'bg-red-400/10 text-red-400 hover:bg-red-400/20 hover:text-red-400 disabled:opacity-100'
: ''
)}
showErrorState &&
'bg-red-400/10 text-red-400 hover:bg-red-400/20 hover:text-red-400 disabled:opacity-100'
]}
{...props}
>
<span class="sr-only">Send</span>
@@ -2,31 +2,27 @@
import { Square } from '@lucide/svelte';
import { Button } from '$lib/components/ui/button';
import {
ChatFormActionAttachmentsDropdown,
ChatFormActionAttachmentsSheet,
ChatFormActionsAdd,
ChatFormActionModels,
ChatFormActionRecord,
ChatFormActionSubmit,
ModelsSelectorDropdown,
ModelsSelectorSheet
ChatFormActionSubmit
} from '$lib/components/app';
import { FileTypeCategory } from '$lib/enums';
import { IsMobile } from '$lib/hooks/is-mobile.svelte';
import { chatStore } from '$lib/stores/chat.svelte';
import { mcpStore } from '$lib/stores/mcp.svelte';
import { modelsStore, modelOptions, selectedModelId } from '$lib/stores/models.svelte';
import { isRouterMode, serverError } from '$lib/stores/server.svelte';
import { config } from '$lib/stores/settings.svelte';
import { activeMessages, conversationsStore } from '$lib/stores/conversations.svelte';
import { conversationsStore } from '$lib/stores/conversations.svelte';
import { getFileTypeCategory } from '$lib/utils';
import { goto } from '$app/navigation';
interface Props {
canSend?: boolean;
canSubmit?: boolean;
class?: string;
disabled?: boolean;
isLoading?: boolean;
isRecording?: boolean;
hasText?: boolean;
showAddButton?: boolean;
showModelSelector?: boolean;
uploadedFiles?: ChatUploadedFile[];
onFileUpload?: () => void;
onMicClick?: () => void;
@@ -38,11 +34,13 @@
let {
canSend = false,
canSubmit = false,
class: className = '',
disabled = false,
isLoading = false,
isRecording = false,
hasText = false,
showAddButton = true,
showModelSelector = true,
uploadedFiles = [],
onFileUpload,
onMicClick,
@@ -53,124 +51,6 @@
}: Props = $props();
let currentConfig = $derived(config());
let isRouter = $derived(isRouterMode());
let isOffline = $derived(!!serverError());
let conversationModel = $derived(
chatStore.getConversationModel(activeMessages() as DatabaseMessage[])
);
let lastSyncedConversationModel: string | null = null;
$effect(() => {
if (conversationModel && conversationModel !== lastSyncedConversationModel) {
lastSyncedConversationModel = conversationModel;
modelsStore.selectModelByName(conversationModel);
} else if (isRouter && !modelsStore.selectedModelId && modelsStore.loadedModelIds.length > 0) {
lastSyncedConversationModel = null;
// auto-select the first loaded model only when nothing is selected yet
const first = modelOptions().find((m) => modelsStore.loadedModelIds.includes(m.model));
if (first) modelsStore.selectModelById(first.id);
}
});
let activeModelId = $derived.by(() => {
const options = modelOptions();
if (!isRouter) {
return options.length > 0 ? options[0].model : null;
}
const selectedId = selectedModelId();
if (selectedId) {
const model = options.find((m) => m.id === selectedId);
if (model) return model.model;
}
if (conversationModel) {
const model = options.find((m) => m.model === conversationModel);
if (model) return model.model;
}
return null;
});
let modelPropsVersion = $state(0); // Used to trigger reactivity after fetch
$effect(() => {
if (activeModelId) {
const cached = modelsStore.getModelProps(activeModelId);
if (!cached) {
modelsStore.fetchModelProps(activeModelId).then(() => {
modelPropsVersion++;
});
}
}
});
let hasAudioModality = $derived.by(() => {
if (activeModelId) {
void modelPropsVersion;
return modelsStore.modelSupportsAudio(activeModelId);
}
return false;
});
let hasVisionModality = $derived.by(() => {
if (activeModelId) {
void modelPropsVersion;
return modelsStore.modelSupportsVision(activeModelId);
}
return false;
});
let hasAudioAttachments = $derived(
uploadedFiles.some((file) => getFileTypeCategory(file.type) === FileTypeCategory.AUDIO)
);
let shouldShowRecordButton = $derived(
hasAudioModality && !hasText && !hasAudioAttachments && currentConfig.autoMicOnEmpty
);
let hasModelSelected = $derived(!isRouter || !!conversationModel || !!selectedModelId());
let isSelectedModelInCache = $derived.by(() => {
if (!isRouter) return true;
if (conversationModel) {
return modelOptions().some((option) => option.model === conversationModel);
}
const currentModelId = selectedModelId();
if (!currentModelId) return false;
return modelOptions().some((option) => option.id === currentModelId);
});
let submitTooltip = $derived.by(() => {
if (!hasModelSelected) {
return 'Please select a model first';
}
if (!isSelectedModelInCache) {
return 'Selected model is not available, please select another';
}
return '';
});
let selectorModelRef: ModelsSelectorDropdown | ModelsSelectorSheet | undefined =
$state(undefined);
let isMobile = new IsMobile();
export function openModelSelector() {
selectorModelRef?.open();
}
let hasMcpPromptsSupport = $derived.by(() => {
const perChatOverrides = conversationsStore.getAllMcpServerOverrides();
@@ -183,25 +63,34 @@
return mcpStore.hasResourcesCapability(perChatOverrides);
});
let hasAudioModality = $state(false);
let hasVisionModality = $state(false);
let hasModelSelected = $state(false);
let isSelectedModelInCache = $state(true);
let submitTooltip = $state('');
let hasAudioAttachments = $derived(
uploadedFiles.some((file) => getFileTypeCategory(file.type) === FileTypeCategory.AUDIO)
);
let shouldShowRecordButton = $derived(
hasAudioModality && !canSubmit && !hasAudioAttachments && currentConfig.autoMicOnEmpty
);
let selectorModelRef: ChatFormActionModels | undefined = $state(undefined);
export function openModelSelector() {
selectorModelRef?.open();
}
</script>
<div class="flex w-full items-center gap-3 {className}" style="container-type: inline-size">
<div class="mr-auto flex items-center gap-2">
{#if isMobile.current}
<ChatFormActionAttachmentsSheet
{disabled}
{hasAudioModality}
{hasVisionModality}
{hasMcpPromptsSupport}
{hasMcpResourcesSupport}
{onFileUpload}
{onSystemPromptClick}
{onMcpPromptClick}
onMcpSettingsClick={() => goto('#/settings/mcp')}
{onMcpResourcesClick}
/>
{:else}
<ChatFormActionAttachmentsDropdown
<div
class="flex w-full items-center gap-3 {className} {showAddButton ? '' : 'justify-end'}"
style="container-type: inline-size"
>
{#if showAddButton}
<div class="mr-auto flex items-center gap-2">
<ChatFormActionsAdd
{disabled}
{hasAudioModality}
{hasVisionModality}
@@ -213,30 +102,24 @@
{onMcpResourcesClick}
onMcpSettingsClick={() => goto('#/settings/mcp')}
/>
{/if}
</div>
</div>
{/if}
<div class="ml-auto flex items-center gap-2">
{#if isMobile.current}
<ModelsSelectorSheet
disabled={disabled || isOffline}
bind:this={selectorModelRef}
currentModel={conversationModel}
forceForegroundText
useGlobalSelection
/>
{:else}
<ModelsSelectorDropdown
disabled={disabled || isOffline}
bind:this={selectorModelRef}
currentModel={conversationModel}
forceForegroundText
useGlobalSelection
/>
{/if}
</div>
{#if showModelSelector}
<ChatFormActionModels
{disabled}
bind:this={selectorModelRef}
bind:hasAudioModality
bind:hasVisionModality
bind:hasModelSelected
bind:isSelectedModelInCache
bind:submitTooltip
forceForegroundText
useGlobalSelection
/>
{/if}
{#if isLoading && !hasText}
{#if isLoading && !canSubmit}
<Button
type="button"
variant="secondary"
@@ -253,10 +136,10 @@
<ChatFormActionRecord {disabled} {hasAudioModality} {isLoading} {isRecording} {onMicClick} />
{:else}
<ChatFormActionSubmit
canSend={canSend && hasModelSelected && isSelectedModelInCache}
canSend={canSend && (showModelSelector ? hasModelSelected && isSelectedModelInCache : true)}
{disabled}
tooltipLabel={submitTooltip}
showErrorState={hasModelSelected && !isSelectedModelInCache}
showErrorState={showModelSelector && hasModelSelected && !isSelectedModelInCache}
/>
{/if}
</div>
@@ -15,6 +15,7 @@
function handleFileSelect(event: Event) {
const input = event.target as HTMLInputElement;
if (input.files) {
onFileSelect?.(Array.from(input.files));
}
@@ -1,31 +0,0 @@
<script lang="ts">
import { browser } from '$app/environment';
import { config } from '$lib/stores/settings.svelte';
interface Props {
class?: string;
show?: boolean;
}
let { class: className = '', show = true }: Props = $props();
let sendOnEnter = $derived(config().sendOnEnter !== false);
let modKey = browser && /Mac|iPhone|iPad|iPod/.test(navigator.platform) ? 'Cmd' : 'Ctrl';
</script>
{#if show}
<div class="mt-6 items-center justify-center {className} hidden md:flex">
{#if sendOnEnter}
<p class="text-xs text-muted-foreground">
Press <kbd class="rounded bg-muted px-1 py-0.5 font-mono text-xs">Enter</kbd> to send,
<kbd class="rounded bg-muted px-1 py-0.5 font-mono text-xs">Shift + Enter</kbd> for new line
</p>
{:else}
<p class="text-xs text-muted-foreground">
Press <kbd class="rounded bg-muted px-1 py-0.5 font-mono text-xs">{modKey} + Enter</kbd> to
send,
<kbd class="rounded bg-muted px-1 py-0.5 font-mono text-xs">Enter</kbd> for new line
</p>
{/if}
</div>
{/if}
@@ -4,7 +4,10 @@
mcpResourceAttachments,
mcpHasResourceAttachments
} from '$lib/stores/mcp-resources.svelte';
import { ChatAttachmentMcpResource, HorizontalScrollCarousel } from '$lib/components/app';
import {
ChatAttachmentsListItemMcpResource,
HorizontalScrollCarousel
} from '$lib/components/app';
interface Props {
class?: string;
@@ -29,11 +32,11 @@
<div class={className}>
<HorizontalScrollCarousel gapSize="2">
{#each attachments as attachment, i (attachment.id)}
<ChatAttachmentMcpResource
<ChatAttachmentsListItemMcpResource
class={i === 0 ? 'ml-3' : ''}
{attachment}
onRemove={handleRemove}
onClick={() => handleResourceClick(attachment.resource.uri)}
onclick={() => handleResourceClick(attachment.resource.uri)}
/>
{/each}
</HorizontalScrollCarousel>

Some files were not shown because too many files have changed in this diff Show More