diff --git a/.github/workflows/check-vendor.yml b/.github/workflows/check-vendor.yml new file mode 100644 index 000000000..7b3016079 --- /dev/null +++ b/.github/workflows/check-vendor.yml @@ -0,0 +1,52 @@ +name: Check vendor + +on: + workflow_dispatch: # allows manual triggering + push: + branches: + - master + paths: [ + 'vendor/**', + 'scripts/sync_vendor.py' + ] + + pull_request: + types: [opened, synchronize, reopened] + paths: [ + 'vendor/**', + 'scripts/sync_vendor.py' + ] + +jobs: + check-vendor: + runs-on: ubuntu-latest + + steps: + - name: Checkout + uses: actions/checkout@v4 + with: + fetch-depth: 0 + + - name: Setup Python + uses: actions/setup-python@v4 + with: + python-version: '3.x' + + - name: Run vendor sync + run: | + set -euo pipefail + python3 scripts/sync_vendor.py + + - name: Check for changes + run: | + set -euo pipefail + # detect modified or untracked files + changed=$(git status --porcelain --untracked-files=all || true) + if [ -n "$changed" ]; then + echo "Vendor sync modified files:" + echo "$changed" | awk '{ print $2 }' | sed '/^$/d' + echo "Failing because vendor files mismatch. Please update scripts/sync_vendor.py" + exit 1 + else + echo "Vendor files are up-to-date." + fi diff --git a/common/download.cpp b/common/download.cpp index debf976b1..eeb32b6a8 100644 --- a/common/download.cpp +++ b/common/download.cpp @@ -467,8 +467,7 @@ std::pair> common_remote_get_content(const std::string & return { res_code, std::move(res_buffer) }; } -#else -#ifdef LLAMA_USE_HTTPLIB +#elif defined(LLAMA_USE_HTTPLIB) static bool is_output_a_tty() { #if defined(_WIN32) @@ -712,29 +711,10 @@ std::pair> common_remote_get_content(const std::string return { res->status, std::move(buf) }; } -#else //no httplib - -bool common_has_curl() { - return false; -} - -static bool common_download_file_single_online(const std::string &, const std::string &, const std::string &) { - LOG_ERR("error: built without CURL, cannot download model from internet\n"); - return false; -} - -std::pair> common_remote_get_content(const std::string & url, const common_remote_params &) { - if (!url.empty()) { - throw std::runtime_error("error: built without CURL, cannot download model from the internet"); - } - - return {}; -} -#endif - - #endif // LLAMA_USE_CURL +#if defined(LLAMA_USE_CURL) || defined(LLAMA_USE_HTTPLIB) + static bool common_download_file_single(const std::string & url, const std::string & path, const std::string & bearer_token, @@ -929,33 +909,6 @@ common_hf_file_res common_get_hf_file(const std::string & hf_repo_with_tag, cons return { hf_repo, ggufFile, mmprojFile }; } -std::vector common_list_cached_models() { - std::vector models; - const std::string cache_dir = fs_get_cache_directory(); - const std::vector files = fs_list_files(cache_dir); - for (const auto & file : files) { - if (string_starts_with(file.name, "manifest=") && string_ends_with(file.name, ".json")) { - common_cached_model_info model_info; - model_info.manifest_path = file.path; - std::string fname = file.name; - string_replace_all(fname, ".json", ""); // remove extension - auto parts = string_split(fname, '='); - if (parts.size() == 4) { - // expect format: manifest==== - model_info.user = parts[1]; - model_info.model = parts[2]; - model_info.tag = parts[3]; - } else { - // invalid format - continue; - } - model_info.size = 0; // TODO: get GGUF size, not manifest size - models.push_back(model_info); - } - } - return models; -} - // // Docker registry functions // @@ -1074,3 +1027,46 @@ std::string common_docker_resolve_model(const std::string & docker) { throw; } } + +#else + +common_hf_file_res common_get_hf_file(const std::string &, const std::string &, bool) { + throw std::runtime_error("download functionality is not enabled in this build"); +} + +bool common_download_model(const common_params_model &, const std::string &, bool) { + throw std::runtime_error("download functionality is not enabled in this build"); +} + +std::string common_docker_resolve_model(const std::string &) { + throw std::runtime_error("download functionality is not enabled in this build"); +} + +#endif // LLAMA_USE_CURL || LLAMA_USE_HTTPLIB + +std::vector common_list_cached_models() { + std::vector models; + const std::string cache_dir = fs_get_cache_directory(); + const std::vector files = fs_list_files(cache_dir); + for (const auto & file : files) { + if (string_starts_with(file.name, "manifest=") && string_ends_with(file.name, ".json")) { + common_cached_model_info model_info; + model_info.manifest_path = file.path; + std::string fname = file.name; + string_replace_all(fname, ".json", ""); // remove extension + auto parts = string_split(fname, '='); + if (parts.size() == 4) { + // expect format: manifest==== + model_info.user = parts[1]; + model_info.model = parts[2]; + model_info.tag = parts[3]; + } else { + // invalid format + continue; + } + model_info.size = 0; // TODO: get GGUF size, not manifest size + models.push_back(model_info); + } + } + return models; +} diff --git a/convert_hf_to_gguf.py b/convert_hf_to_gguf.py index 13448fd68..cc77a3db2 100755 --- a/convert_hf_to_gguf.py +++ b/convert_hf_to_gguf.py @@ -7354,6 +7354,7 @@ class PLMModel(TextModel): @ModelBase.register("T5ForConditionalGeneration") @ModelBase.register("MT5ForConditionalGeneration") @ModelBase.register("UMT5ForConditionalGeneration") +@ModelBase.register("UMT5Model") class T5Model(TextModel): model_arch = gguf.MODEL_ARCH.T5 diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index 6e09bbbf5..c7843dae9 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -4140,6 +4140,13 @@ void ggml_cpu_fp16_to_fp32(const ggml_fp16_t * x, float * y, int64_t n) { __m128 y_vec = _mm_cvtph_ps(x_vec); _mm_storeu_ps(y + i, y_vec); } +#elif defined(__riscv_zvfh) + for (int vl; i < n; i += vl) { + vl = __riscv_vsetvl_e16m1(n - i); + vfloat16m1_t vx = __riscv_vle16_v_f16m1((_Float16 *)&x[i], vl); + vfloat32m2_t vy = __riscv_vfwcvt_f_f_v_f32m2(vx, vl); + __riscv_vse32_v_f32m2(&y[i], vy, vl); + } #endif for (; i < n; ++i) { diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index 7c42fb78b..9f1e5f8d6 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -7,8 +7,9 @@ #include "unary-ops.h" #include "vec.h" -#include +#include #include +#include // ggml_compute_forward_dup @@ -5503,7 +5504,28 @@ static void ggml_mrope_cache_init( } } -static void ggml_compute_forward_rope_f32( + +template +static void rotate_pairs(const int64_t n, const int64_t n_offset, const float * cache, const T * src_data, T * dst_data, const int scale = 2) { + for (int64_t i0 = 0; i0 < n; i0 += 2) { + const int64_t ic = i0/scale; // hack for GGML_ROPE_TYPE_NORMAL, where we need ic = i0; for all other cases, ic = i0/2 + + const float cos_theta = cache[i0 + 0]; + const float sin_theta = cache[i0 + 1]; + + const T * const src = src_data + ic; + T * dst = dst_data + ic; + + const float x0 = type_conversion_table::to_f32(src[0]); + const float x1 = type_conversion_table::to_f32(src[n_offset]); + + dst[0] = type_conversion_table::from_f32(x0*cos_theta - x1*sin_theta); + dst[n_offset] = type_conversion_table::from_f32(x0*sin_theta + x1*cos_theta); + } +} + +template //float or ggml_fp16_t +static void ggml_compute_forward_rope_flt( const ggml_compute_params * params, ggml_tensor * dst, const bool forward) { @@ -5512,6 +5534,9 @@ static void ggml_compute_forward_rope_f32( const ggml_tensor * src1 = dst->src[1]; const ggml_tensor * src2 = dst->src[2]; + GGML_ASSERT(src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16); + GGML_ASSERT(src1->type == GGML_TYPE_I32); + float freq_base, freq_scale, ext_factor, attn_factor, beta_fast, beta_slow; int sections[4]; @@ -5534,7 +5559,8 @@ static void ggml_compute_forward_rope_f32( //printf("ne0: %d, ne1: %d, ne2: %d, ne3: %d\n", ne0, ne1, ne2, ne3); //printf("n_past = %d, ne2 = %d\n", n_past, ne2); - GGML_ASSERT(nb00 == sizeof(float)); + GGML_ASSERT(nb0 == nb00); + GGML_ASSERT(nb0 == sizeof(T)); const int ith = params->ith; const int nth = params->nth; @@ -5559,12 +5585,11 @@ static void ggml_compute_forward_rope_f32( float corr_dims[2]; ggml_rope_yarn_corr_dims(n_dims, n_ctx_orig, freq_base, beta_fast, beta_slow, corr_dims); - const bool is_neox = mode & GGML_ROPE_TYPE_NEOX; - const bool is_mrope = mode & GGML_ROPE_TYPE_MROPE; // ggml_rope_multi, multimodal rotary position embedding const bool is_imrope = mode == GGML_ROPE_TYPE_IMROPE; // qwen3vl apply interleaved mrope + const bool mrope_used = mode & GGML_ROPE_TYPE_MROPE; // ggml_rope_multi, note: also true for vision (24 & 8 == true) and for imrope const bool is_vision = mode == GGML_ROPE_TYPE_VISION; - if (is_mrope) { + if (mrope_used) { GGML_ASSERT(sections[0] > 0 || sections[1] > 0 || sections[2] > 0); } @@ -5590,7 +5615,7 @@ static void ggml_compute_forward_rope_f32( for (int64_t i2 = 0; i2 < ne2; i2++) { // seq-len float * cache = (float *) params->wdata + (ne0 + CACHE_LINE_SIZE_F32)*ith; - if (!is_mrope) { + if (!mrope_used) { const int64_t p = pos[i2]; ggml_rope_cache_init(p, freq_scale, freq_factors, corr_dims, ne0, ext_factor, attn_factor, cache, sin_sign, theta_scale); } @@ -5608,269 +5633,36 @@ static void ggml_compute_forward_rope_f32( if (ir++ < ir0) continue; if (ir > ir1) break; - if (is_neox || is_mrope) { - if (is_vision){ - for (int64_t i0 = 0; i0 < n_dims; i0 += 2) { - const int64_t ic = i0/2; + T * src = (T *)((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01); + T * dst_data = (T *)((char *) dst->data + i3*nb3 + i2*nb2 + i1*nb1); - const float cos_theta = cache[i0 + 0]; - const float sin_theta = cache[i0 + 1]; - - const float * const src = (float *)((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + ic*nb00); - float * dst_data = (float *)((char *) dst->data + i3*nb3 + i2*nb2 + i1*nb1 + ic*nb0); - - const float x0 = src[0]; - const float x1 = src[n_dims]; - - dst_data[0] = x0*cos_theta - x1*sin_theta; - dst_data[n_dims] = x0*sin_theta + x1*cos_theta; - } - } else { - for (int64_t i0 = 0; i0 < n_dims; i0 += 2) { - const int64_t ic = i0/2; - - const float cos_theta = cache[i0 + 0]; - const float sin_theta = cache[i0 + 1]; - - const float * const src = (float *)((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + ic*nb00); - float * dst_data = (float *)((char *) dst->data + i3*nb3 + i2*nb2 + i1*nb1 + ic*nb0); - - const float x0 = src[0]; - const float x1 = src[n_dims/2]; - - dst_data[0] = x0*cos_theta - x1*sin_theta; - dst_data[n_dims/2] = x0*sin_theta + x1*cos_theta; - } - } - } else { - for (int64_t i0 = 0; i0 < n_dims; i0 += 2) { - const float cos_theta = cache[i0 + 0]; - const float sin_theta = cache[i0 + 1]; - - const float * const src = (float *)((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00); - float * dst_data = (float *)((char *) dst->data + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0); - - const float x0 = src[0]; - const float x1 = src[1]; - - dst_data[0] = x0*cos_theta - x1*sin_theta; - dst_data[1] = x0*sin_theta + x1*cos_theta; - } + switch (mode) { + case GGML_ROPE_TYPE_NORMAL: + rotate_pairs(n_dims, 1, cache, src, dst_data, 1); + break; + case GGML_ROPE_TYPE_NEOX: + case GGML_ROPE_TYPE_MROPE: + case GGML_ROPE_TYPE_IMROPE: + rotate_pairs(n_dims, n_dims/2, cache, src, dst_data); + break; + case GGML_ROPE_TYPE_VISION: + rotate_pairs(ne0, n_dims, cache, src, dst_data); + break; + default: + GGML_ABORT("rope type not supported"); } - if (is_vision) { - for (int64_t i0 = n_dims; i0 < ne0; i0 += 2) { - const int64_t ic = i0/2; - - const float cos_theta = cache[i0 + 0]; - const float sin_theta = cache[i0 + 1]; - - const float * const src = (float *)((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + ic*nb00); - float * dst_data = (float *)((char *) dst->data + i3*nb3 + i2*nb2 + i1*nb1 + ic*nb0); - - const float x0 = src[0]; - const float x1 = src[n_dims]; - - dst_data[0] = x0*cos_theta - x1*sin_theta; - dst_data[n_dims] = x0*sin_theta + x1*cos_theta; - } - } else { + if (!is_vision) { // fill the remain channels with data from src tensor for (int64_t i0 = n_dims; i0 < ne0; i0 += 2) { - const float * const src = (float *)((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00); - float * dst_data = (float *)((char *) dst->data + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0); + const T * const src = (T *)((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00); + T * dst_data = (T *)((char *) dst->data + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0); dst_data[0] = src[0]; dst_data[1] = src[1]; } } - } - } - } -} - -// TODO: deduplicate f16/f32 code -static void ggml_compute_forward_rope_f16( - const ggml_compute_params * params, - ggml_tensor * dst, - const bool forward) { - - const ggml_tensor * src0 = dst->src[0]; - const ggml_tensor * src1 = dst->src[1]; - const ggml_tensor * src2 = dst->src[2]; - - float freq_base, freq_scale, ext_factor, attn_factor, beta_fast, beta_slow; - int sections[4]; - - //const int n_past = ((int32_t *) dst->op_params)[0]; - const int n_dims = ((int32_t *) dst->op_params)[1]; - const int mode = ((int32_t *) dst->op_params)[2]; - //const int n_ctx = ((int32_t *) dst->op_params)[3]; - const int n_ctx_orig = ((int32_t *) dst->op_params)[4]; - memcpy(&freq_base, (int32_t *) dst->op_params + 5, sizeof(float)); - memcpy(&freq_scale, (int32_t *) dst->op_params + 6, sizeof(float)); - memcpy(&ext_factor, (int32_t *) dst->op_params + 7, sizeof(float)); - memcpy(&attn_factor, (int32_t *) dst->op_params + 8, sizeof(float)); - memcpy(&beta_fast, (int32_t *) dst->op_params + 9, sizeof(float)); - memcpy(&beta_slow, (int32_t *) dst->op_params + 10, sizeof(float)); - memcpy(§ions, (int32_t *) dst->op_params + 11, sizeof(int)*4); - - - GGML_TENSOR_UNARY_OP_LOCALS - - //printf("ne0: %d, ne1: %d, ne2: %d, ne3: %d\n", ne0, ne1, ne2, ne3); - //printf("n_past = %d, ne2 = %d\n", n_past, ne2); - - GGML_ASSERT(nb0 == sizeof(ggml_fp16_t)); - - const int ith = params->ith; - const int nth = params->nth; - - const int nr = ggml_nrows(dst); - - GGML_ASSERT(n_dims <= ne0); - GGML_ASSERT(n_dims % 2 == 0); - - // rows per thread - const int dr = (nr + nth - 1)/nth; - - // row range for this thread - const int ir0 = dr*ith; - const int ir1 = MIN(ir0 + dr, nr); - - // row index used to determine which thread to use - int ir = 0; - - const float theta_scale = powf(freq_base, -2.0f/n_dims); - - float corr_dims[2]; - ggml_rope_yarn_corr_dims(n_dims, n_ctx_orig, freq_base, beta_fast, beta_slow, corr_dims); - - const bool is_neox = mode & GGML_ROPE_TYPE_NEOX; - const bool is_mrope = mode & GGML_ROPE_TYPE_MROPE; - const bool is_imrope = mode == GGML_ROPE_TYPE_IMROPE; - const bool is_vision = mode == GGML_ROPE_TYPE_VISION; - - if (is_mrope) { - GGML_ASSERT(sections[0] > 0 || sections[1] > 0 || sections[2] > 0); - } - - if (is_vision) { - GGML_ASSERT(n_dims == ne0/2); - } - - const float * freq_factors = NULL; - if (src2 != NULL) { - GGML_ASSERT(src2->type == GGML_TYPE_F32); - GGML_ASSERT(src2->ne[0] >= n_dims / 2); - freq_factors = (const float *) src2->data; - } - - // backward process uses inverse rotation by cos and sin. - // cos and sin build a rotation matrix, where the inverse is the transpose. - // this essentially just switches the sign of sin. - const float sin_sign = forward ? 1.0f : -1.0f; - - const int32_t * pos = (const int32_t *) src1->data; - - for (int64_t i3 = 0; i3 < ne3; i3++) { - for (int64_t i2 = 0; i2 < ne2; i2++) { - - float * cache = (float *) params->wdata + (ne0 + CACHE_LINE_SIZE_F32)*ith; - if (!is_mrope) { - const int64_t p = pos[i2]; - ggml_rope_cache_init(p, freq_scale, freq_factors, corr_dims, ne0, ext_factor, attn_factor, cache, sin_sign, theta_scale); - } - else { - const int64_t p_t = pos[i2]; - const int64_t p_h = pos[i2 + ne2]; - const int64_t p_w = pos[i2 + ne2 * 2]; - const int64_t p_e = pos[i2 + ne2 * 3]; - ggml_mrope_cache_init( - p_t, p_h, p_w, p_e, sections, is_imrope, is_vision, - freq_scale, freq_factors, corr_dims, ne0, ext_factor, attn_factor, cache, sin_sign, theta_scale); - } - - for (int64_t i1 = 0; i1 < ne1; i1++) { - if (ir++ < ir0) continue; - if (ir > ir1) break; - - if (is_neox || is_mrope) { - if (is_vision) { - for (int64_t i0 = 0; i0 < n_dims; i0 += 2) { - const int64_t ic = i0/2; - - const float cos_theta = cache[i0 + 0]; - const float sin_theta = cache[i0 + 1]; - - const ggml_fp16_t * const src = (ggml_fp16_t *)((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + ic*nb00); - ggml_fp16_t * dst_data = (ggml_fp16_t *)((char *) dst->data + i3*nb3 + i2*nb2 + i1*nb1 + ic*nb0); - - const float x0 = GGML_CPU_FP16_TO_FP32(src[0]); - const float x1 = GGML_CPU_FP16_TO_FP32(src[n_dims]); - - dst_data[0] = GGML_CPU_FP32_TO_FP16(x0*cos_theta - x1*sin_theta); - dst_data[n_dims] = GGML_CPU_FP32_TO_FP16(x0*sin_theta + x1*cos_theta); - } - } else { - for (int64_t i0 = 0; i0 < n_dims; i0 += 2) { - const int64_t ic = i0/2; - - const float cos_theta = cache[i0 + 0]; - const float sin_theta = cache[i0 + 1]; - - const ggml_fp16_t * const src = (ggml_fp16_t *)((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + ic*nb00); - ggml_fp16_t * dst_data = (ggml_fp16_t *)((char *) dst->data + i3*nb3 + i2*nb2 + i1*nb1 + ic*nb0); - - const float x0 = GGML_CPU_FP16_TO_FP32(src[0]); - const float x1 = GGML_CPU_FP16_TO_FP32(src[n_dims/2]); - - dst_data[0] = GGML_CPU_FP32_TO_FP16(x0*cos_theta - x1*sin_theta); - dst_data[n_dims/2] = GGML_CPU_FP32_TO_FP16(x0*sin_theta + x1*cos_theta); - } - } - } else { - for (int64_t i0 = 0; i0 < n_dims; i0 += 2) { - const float cos_theta = cache[i0 + 0]; - const float sin_theta = cache[i0 + 1]; - - const ggml_fp16_t * const src = (ggml_fp16_t *)((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00); - ggml_fp16_t * dst_data = (ggml_fp16_t *)((char *) dst->data + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0); - - const float x0 = GGML_CPU_FP16_TO_FP32(src[0]); - const float x1 = GGML_CPU_FP16_TO_FP32(src[1]); - - dst_data[0] = GGML_CPU_FP32_TO_FP16(x0*cos_theta - x1*sin_theta); - dst_data[1] = GGML_CPU_FP32_TO_FP16(x0*sin_theta + x1*cos_theta); - } - } - - if (is_vision) { - for (int64_t i0 = n_dims; i0 < ne0; i0 += 2) { - const int64_t ic = i0/2; - - const float cos_theta = cache[i0 + 0]; - const float sin_theta = cache[i0 + 1]; - - const ggml_fp16_t * const src = (ggml_fp16_t *)((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + ic*nb00); - ggml_fp16_t * dst_data = (ggml_fp16_t *)((char *) dst->data + i3*nb3 + i2*nb2 + i1*nb1 + ic*nb0); - - const float x0 = GGML_CPU_FP16_TO_FP32(src[0]); - const float x1 = GGML_CPU_FP16_TO_FP32(src[n_dims]); - - dst_data[0] = GGML_CPU_FP32_TO_FP16(x0*cos_theta - x1*sin_theta); - dst_data[n_dims] = GGML_CPU_FP32_TO_FP16(x0*sin_theta + x1*cos_theta); - } - } else { - for (int64_t i0 = n_dims; i0 < ne0; i0 += 2) { - const ggml_fp16_t * const src = (ggml_fp16_t *)((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00); - ggml_fp16_t * dst_data = (ggml_fp16_t *)((char *) dst->data + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0); - - dst_data[0] = src[0]; - dst_data[1] = src[1]; - } - } - } + } //attn-heads } } } @@ -5884,11 +5676,11 @@ void ggml_compute_forward_rope( switch (src0->type) { case GGML_TYPE_F16: { - ggml_compute_forward_rope_f16(params, dst, true); + ggml_compute_forward_rope_flt(params, dst, true); } break; case GGML_TYPE_F32: { - ggml_compute_forward_rope_f32(params, dst, true); + ggml_compute_forward_rope_flt(params, dst, true); } break; default: { @@ -5908,11 +5700,11 @@ void ggml_compute_forward_rope_back( switch (src0->type) { case GGML_TYPE_F16: { - ggml_compute_forward_rope_f16(params, dst, false); + ggml_compute_forward_rope_flt(params, dst, false); } break; case GGML_TYPE_F32: { - ggml_compute_forward_rope_f32(params, dst, false); + ggml_compute_forward_rope_flt(params, dst, false); } break; default: { @@ -7891,24 +7683,24 @@ static void ggml_compute_forward_argsort_f32( ggml_sort_order order = (ggml_sort_order) ggml_get_op_params_i32(dst, 0); for (int64_t i = ith; i < nr; i += nth) { - int32_t * dst_data = (int32_t *)((char *) dst->data + i*nb1); const float * src_data = (float *)((char *) src0->data + i*nb01); + int32_t * dst_data = (int32_t *)((char *) dst->data + i*nb1); + for (int64_t j = 0; j < ne0; j++) { dst_data[j] = j; } - // C doesn't have a functional sort, so we do a bubble sort instead - for (int64_t j = 0; j < ne0; j++) { - for (int64_t k = j + 1; k < ne0; k++) { - if ((order == GGML_SORT_ORDER_ASC && src_data[dst_data[j]] > src_data[dst_data[k]]) || - (order == GGML_SORT_ORDER_DESC && src_data[dst_data[j]] < src_data[dst_data[k]])) { - int32_t tmp = dst_data[j]; - dst_data[j] = dst_data[k]; - dst_data[k] = tmp; - } - } + std::function cmp; + + // note: this might be causing memory allocations? ideally should be avoided if it's the case + switch (order) { + case GGML_SORT_ORDER_ASC: cmp = [src_data](int32_t a, int32_t b) { return src_data[a] < src_data[b]; }; break; + case GGML_SORT_ORDER_DESC: cmp = [src_data](int32_t a, int32_t b) { return src_data[a] > src_data[b]; }; break; + default: GGML_ABORT("invalid sort order"); } + + std::sort(dst_data, dst_data + ne0, cmp); } } diff --git a/ggml/src/ggml-cpu/repack.cpp b/ggml/src/ggml-cpu/repack.cpp index 1e58a67f5..01944eaab 100644 --- a/ggml/src/ggml-cpu/repack.cpp +++ b/ggml/src/ggml-cpu/repack.cpp @@ -1474,29 +1474,52 @@ template src[0]; const ggml_tensor * src1 = op->src[1]; ggml_tensor * dst = op; GGML_TENSOR_BINARY_OP_LOCALS - const void * src1_wdata = params->wdata; const size_t src1_col_stride = ggml_row_size(PARAM_TYPE, ne10); + GGML_ASSERT(ne03 == 1 && ne13 == 1); + GGML_ASSERT(ne12 % ne02 == 0); + const int64_t r2 = ne12 / ne02; + + const int64_t i12 = src1_start / ne1; + const int64_t i11 = src1_start - i12 * ne1; + + // Determine batch index + const int64_t i02 = i12 / r2; + + const int64_t i1 = i11; + const int64_t i2 = i12; + + const char * src0_ptr = (const char *) src0->data + i02 * nb02; + const char * src1_ptr = (const char *) params->wdata + (i11 + i12 * ne11) * src1_col_stride; + char * dst_ptr = ((char *) dst->data + (i1 * nb1 + i2 * nb2)); + + const int64_t nrows = src1_end - src1_start; + const int64_t ncols = src0_end - src0_start; + + GGML_ASSERT(src1_ptr + src1_col_stride * nrows <= (const char *) params->wdata + params->wsize); + // If there are more than three rows in src1, use gemm; otherwise, use gemv. - if (ne11 > 3) { - gemm(ne00, - (float *) ((char *) dst->data) + src0_start, ne01, - (const char *) src0->data + src0_start * nb01, - (const char *) src1_wdata, ne11 - ne11 % 4, src0_end - src0_start); + if (nrows > 3) { + gemm(ne00, (float *) (dst_ptr) + src0_start, nb1 / nb0, + src0_ptr + src0_start * nb01, src1_ptr, + nrows - (nrows % 4), ncols); } - for (int iter = ne11 - ne11 % 4; iter < ne11; iter++) { - gemv(ne00, - (float *) ((char *) dst->data + (iter * nb1)) + src0_start, ne01, - (const char *) src0->data + src0_start * nb01, - (const char *) src1_wdata + (src1_col_stride * iter), 1, - src0_end - src0_start); + for (int iter = nrows - (nrows % 4); iter < nrows; iter++) { + gemv(ne00, (float *) (dst_ptr + (iter * nb1)) + src0_start, + ne01, src0_ptr + src0_start * nb01, + src1_ptr + (src1_col_stride * iter), 1 /* nrows */, ncols); } } @@ -1521,6 +1544,12 @@ template type == GGML_TYPE_F32); GGML_ASSERT(ggml_n_dims(op->src[0]) == 2); @@ -1528,47 +1557,60 @@ template (params->wdata); const size_t nbw1 = ggml_row_size(PARAM_TYPE, ne10); + const size_t nbw2 = nbw1 * ne11; - assert(params->wsize >= nbw1 * ne11); + assert(params->wsize >= nbw2 * ne12); const ggml_from_float_t from_float = ggml_get_type_traits_cpu(PARAM_TYPE)->from_float; - int64_t i11_processed = 0; - for (int64_t i11 = ith * 4; i11 < ne11 - ne11 % 4; i11 += nth * 4) { - ggml_quantize_mat_t((float *) ((char *) src1->data + i11 * nb11), (void *) (wdata + i11 * nbw1), 4, ne10); - } + for (int64_t i12 = 0; i12 < ne12; i12++) { + char * data_ptr = (char *) src1->data + i12 * nb12; + char * wdata_ptr = wdata + i12 * nbw2; - i11_processed = ne11 - ne11 % 4; - for (int64_t i11 = i11_processed + ith; i11 < ne11; i11 += nth) { - from_float((float *) ((char *) src1->data + i11 * nb11), (void *) (wdata + i11 * nbw1), ne10); + for (int64_t i11 = ith * 4; i11 < ne11 - ne11 % 4; i11 += nth * 4) { + ggml_quantize_mat_t((float *) (data_ptr + i11 * nb11), + (void *) (wdata_ptr + i11 * nbw1), 4, ne10); + } + + const int64_t i11_processed = ne11 - ne11 % 4; + for (int64_t i11 = i11_processed + ith; i11 < ne11; i11 += nth) { + from_float((float *) (data_ptr + i11 * nb11), (void *) (wdata_ptr + i11 * nbw1), ne10); + } } // disable for NUMA const bool disable_chunking = ggml_is_numa(); // 4x chunks per thread - int64_t nr = ggml_nrows(op->src[0]); - int nth_scaled = nth * 4; - int64_t chunk_size = (nr + nth_scaled - 1) / nth_scaled; - int64_t nchunk = (nr + chunk_size - 1) / chunk_size; + const int64_t nr0 = ggml_nrows(op->src[0]); + const int64_t nr1 = ne1 * ne2 * ne3; + + int nth_scaled = nth * 4; + int64_t chunk_size0 = (nr0 + nth_scaled - 1) / nth_scaled; + // avoid too small chunks for narrow src1 + int64_t chunk_size1 = MAX(16, (nr1 + nth - 1) / nth); + int64_t nchunk0 = (nr0 + chunk_size0 - 1) / chunk_size0; + int64_t nchunk1 = (nr1 + chunk_size1 - 1) / chunk_size1; // Ensure minimum chunk size to avoid alignment issues with high thread counts // Minimum chunk size should be at least NB_COLS to prevent overlapping chunks after alignment const int64_t min_chunk_size = NB_COLS; - if (nchunk > 0 && (nr / nchunk) < min_chunk_size && nr >= min_chunk_size) { - nchunk = (nr + min_chunk_size - 1) / min_chunk_size; + if (nchunk0 > 0 && (nr0 / nchunk0) < min_chunk_size && nr0 >= min_chunk_size) { + nchunk0 = (nr0 + min_chunk_size - 1) / min_chunk_size; } - if (nth == 1 || nchunk < nth || disable_chunking) { - nchunk = nth; + if (nth == 1 || nchunk0 * nchunk1 < nth || disable_chunking) { + nchunk0 = nr0 > nr1 ? nth : 1; + nchunk1 = nr0 > nr1 ? 1 : nth; } + const int64_t dr0 = (nr0 + nchunk0 - 1) / nchunk0; + const int64_t dr1 = (nr1 + nchunk1 - 1) / nchunk1; + // Ensure nchunk doesn't exceed the number of rows divided by minimum chunk size // This prevents creating too many tiny chunks that could overlap after alignment - const int64_t max_nchunk = (nr + min_chunk_size - 1) / min_chunk_size; - if (nchunk > max_nchunk) { - nchunk = max_nchunk; - } + const int64_t max_nchunk = (nr0 + min_chunk_size - 1) / min_chunk_size; + nchunk0 = MIN(nchunk0, max_nchunk); if (ith == 0) { // Every thread starts at ith, so the first unprocessed chunk is nth. This save a bit of coordination right at the start. @@ -1580,23 +1622,29 @@ template ne01) { - src0_end = ne01; - } + src0_end = (src0_end % NB_COLS) ? src0_end + NB_COLS - (src0_end % NB_COLS) : src0_end; + src0_end = MIN(src0_end, ne01); + // Make sure current plane is the last one before exiting if (src0_start >= src0_end) { - break; + current_chunk = ggml_threadpool_chunk_add(params->threadpool, 1); + continue; } - forward_mul_mat_one_chunk(params, dst, src0_start, src0_end); + forward_mul_mat_one_chunk(params, dst, src0_start, src0_end, src1_start, src1_end); current_chunk = ggml_threadpool_chunk_add(params->threadpool, 1); } diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 48a25c1b6..31f30d538 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -591,6 +591,12 @@ static __device__ __forceinline__ void ggml_cuda_mad(half2 & acc, const half2 v, // If dst and src point at different address spaces then they are guaranteed to not be aliased. template static __device__ __forceinline__ void ggml_cuda_memcpy_1(void * __restrict__ dst, const void * __restrict__ src) { + static_assert( + nbytes <= ggml_cuda_get_max_cpy_bytes() || alignment == 0, + "You are misusing the alignment parameter for ggml_cuda_memcpy_1. " + "The intent is for the parameter is only as a workaround if either one of the pointers is not properly aligned. " + "If you use it to do more bytes per copy than ggml_cuda_max_cpy_bytes() the reads and writes may not be coalesced. " + "Call ggml_cuda_memcpy_1 in a loop instead."); if constexpr (alignment != 0) { static_assert(nbytes % alignment == 0, "bad alignment"); } diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 941d34da0..490ae73e4 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3005,6 +3005,36 @@ static void update_cuda_graph_executable(ggml_backend_cuda_context * cuda_ctx) { } #endif +static bool ggml_cuda_should_fuse_rope_set_rows(const ggml_tensor * rope, + const ggml_tensor * view, + const ggml_tensor * set_rows) { + // ne3 not tested + if (rope->src[0]->ne[3] != 1) { + return false; + } + + if (set_rows->type != GGML_TYPE_F32 && set_rows->type != GGML_TYPE_F16) { + return false; + } + + if (set_rows->src[1]->type != GGML_TYPE_I64) { + return false; + } + + // The view should flatten two dims of rope into one dim + if (!ggml_is_contiguous(view) || view->ne[0] != rope->ne[0] * rope->ne[1]) { + return false; + } + + // Only norm/neox shaders have the fusion code + const int mode = ((const int32_t *) rope->op_params)[2]; + if (mode != GGML_ROPE_TYPE_NORMAL && mode != GGML_ROPE_TYPE_NEOX) { + return false; + } + + return true; +} + static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, int node_idx, std::initializer_list ops, std::initializer_list unary_ops) { #ifndef NDEBUG const size_t num_unary = std::count(ops.begin(), ops.end(), GGML_OP_UNARY); @@ -3080,6 +3110,16 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, int node_idx, } } + if (ops.size() == 3 && ggml_can_fuse_subgraph(cgraph, node_idx, ops, { node_idx + 2 })) { + const ggml_tensor * rope = cgraph->nodes[node_idx]; + const ggml_tensor * view = cgraph->nodes[node_idx + 1]; + const ggml_tensor * set_rows = cgraph->nodes[node_idx + 2]; + + if (ggml_cuda_should_fuse_rope_set_rows(rope, view, set_rows)) { + return true; + } + } + if (!ggml_can_fuse(cgraph, node_idx, ops)) { return false; } @@ -3209,6 +3249,15 @@ static void evaluate_and_capture_cuda_graph(ggml_backend_cuda_context * cuda_ctx continue; } + if (ggml_cuda_can_fuse(cgraph, i, { GGML_OP_ROPE, GGML_OP_VIEW, GGML_OP_SET_ROWS }, {})) { + ggml_tensor * rope = cgraph->nodes[i]; + ggml_tensor * set_rows = cgraph->nodes[i + 2]; + + ggml_cuda_op_rope_fused(*cuda_ctx, rope, set_rows); + i += 2; + continue; + } + if (node->op == GGML_OP_ADD) { int n_fuse = 0; ggml_op ops[8]; diff --git a/ggml/src/ggml-cuda/rope.cu b/ggml/src/ggml-cuda/rope.cu index 78ed7f519..88ed79111 100644 --- a/ggml/src/ggml-cuda/rope.cu +++ b/ggml/src/ggml-cuda/rope.cu @@ -1,3 +1,6 @@ +#include "convert.cuh" +#include "ggml-cuda/common.cuh" +#include "ggml.h" #include "rope.cuh" struct rope_corr_dims { @@ -37,11 +40,23 @@ static __device__ void rope_yarn( } } -template -static __global__ void rope_norm( - const T * x, T * dst, const int ne0, const int ne1, const int s1, const int s2, const int n_dims, - const int32_t * pos, const float freq_scale, const float ext_factor, const float attn_factor, - const rope_corr_dims corr_dims, const float theta_scale, const float * freq_factors) { +template +static __global__ void rope_norm(const T * x, + D * dst, + const int ne0, + const int ne1, + const int s1, + const int s2, + const int n_dims, + const int32_t * pos, + const float freq_scale, + const float ext_factor, + const float attn_factor, + const rope_corr_dims corr_dims, + const float theta_scale, + const float * freq_factors, + const int64_t * row_indices, + const int set_rows_stride) { const int i0 = 2*(blockDim.y*blockIdx.y + threadIdx.y); if (i0 >= ne0) { @@ -53,13 +68,27 @@ static __global__ void rope_norm( const int row_x = row_dst % ne1; const int channel_x = row_dst / ne1; - const int idst = row_dst*ne0 + i0; + int idst = row_dst * ne0 + i0; const int ix = channel_x*s2 + row_x*s1 + i0; - if (i0 >= n_dims) { - dst[idst + 0] = x[ix + 0]; - dst[idst + 1] = x[ix + 1]; + // Fusion optimization: ROPE + VIEW + SET_ROWS. + // The rope output is viewed as a 1D tensor and offset based on a row index in row_indices. + if (set_rows_stride != 0) { + idst = row_x * ne0 + i0; + idst += row_indices[channel_x] * set_rows_stride; + } + const auto & store_coaelsced = [&](float x0, float x1) { + if constexpr (std::is_same_v) { + float2 v = make_float2(x0, x1); + ggml_cuda_memcpy_1<8>(dst + idst, &v); + } else if constexpr (std::is_same_v) { + half2 v = make_half2(x0, x1); + ggml_cuda_memcpy_1<4>(dst + idst, &v); + } + }; + if (i0 >= n_dims) { + store_coaelsced(x[ix + 0], x[ix + 1]); return; } @@ -75,15 +104,26 @@ static __global__ void rope_norm( const float x0 = x[ix + 0]; const float x1 = x[ix + 1]; - dst[idst + 0] = x0*cos_theta - x1*sin_theta; - dst[idst + 1] = x0*sin_theta + x1*cos_theta; + store_coaelsced(x0 * cos_theta - x1 * sin_theta, x0 * sin_theta + x1 * cos_theta); } -template -static __global__ void rope_neox( - const T * x, T * dst, const int ne0, const int ne1, const int s1, const int s2, const int n_dims, - const int32_t * pos, const float freq_scale, const float ext_factor, const float attn_factor, - const rope_corr_dims corr_dims, const float theta_scale, const float * freq_factors) { +template +static __global__ void rope_neox(const T * x, + D * dst, + const int ne0, + const int ne1, + const int s1, + const int s2, + const int n_dims, + const int32_t * pos, + const float freq_scale, + const float ext_factor, + const float attn_factor, + const rope_corr_dims corr_dims, + const float theta_scale, + const float * freq_factors, + const int64_t * row_indices, + const int set_rows_stride) { const int i0 = 2*(blockDim.y*blockIdx.y + threadIdx.y); if (i0 >= ne0) { @@ -95,12 +135,19 @@ static __global__ void rope_neox( const int row_x = row_dst % ne1; const int channel_x = row_dst / ne1; - const int idst = row_dst*ne0 + i0/2; + int idst = row_dst * ne0 + i0 / 2; const int ix = channel_x*s2 + row_x*s1 + i0/2; + // Fusion optimization: ROPE + VIEW + SET_ROWS. + // The rope output is viewed as a 1D tensor and offset based on a row index in row_indices. + if (set_rows_stride != 0) { + idst = row_x * ne0 + i0 / 2; + idst += row_indices[channel_x] * set_rows_stride; + } + if (i0 >= n_dims) { - dst[idst + i0/2 + 0] = x[ix + i0/2 + 0]; - dst[idst + i0/2 + 1] = x[ix + i0/2 + 1]; + dst[idst + i0 / 2 + 0] = ggml_cuda_cast(x[ix + i0 / 2 + 0]); + dst[idst + i0 / 2 + 1] = ggml_cuda_cast(x[ix + i0 / 2 + 1]); return; } @@ -117,8 +164,8 @@ static __global__ void rope_neox( const float x0 = x[ix + 0]; const float x1 = x[ix + n_dims/2]; - dst[idst + 0] = x0*cos_theta - x1*sin_theta; - dst[idst + n_dims/2] = x0*sin_theta + x1*cos_theta; + dst[idst + 0] = ggml_cuda_cast(x0 * cos_theta - x1 * sin_theta); + dst[idst + n_dims / 2] = ggml_cuda_cast(x0 * sin_theta + x1 * cos_theta); } template @@ -238,11 +285,25 @@ static __global__ void rope_vision( dst[idst + n_dims] = x0*sin_theta + x1*cos_theta; } -template -static void rope_norm_cuda( - const T * x, T * dst, const int ne0, const int ne1, const int s1, const int s2, const int n_dims, const int nr, - const int32_t * pos, const float freq_scale, const float freq_base, const float ext_factor, const float attn_factor, - const rope_corr_dims corr_dims, const float * freq_factors, cudaStream_t stream) { +template +static void rope_norm_cuda(const T * x, + D * dst, + const int ne0, + const int ne1, + const int s1, + const int s2, + const int n_dims, + const int nr, + const int32_t * pos, + const float freq_scale, + const float freq_base, + const float ext_factor, + const float attn_factor, + const rope_corr_dims corr_dims, + const float * freq_factors, + const int64_t * row_indices, + const int set_rows_stride, + cudaStream_t stream) { GGML_ASSERT(ne0 % 2 == 0); const dim3 block_dims(1, CUDA_ROPE_BLOCK_SIZE, 1); const int n_blocks_x = (ne0 + 2*CUDA_ROPE_BLOCK_SIZE - 1) / (2*CUDA_ROPE_BLOCK_SIZE); @@ -252,20 +313,34 @@ static void rope_norm_cuda( if (freq_factors == nullptr) { rope_norm<<>>( - x, dst, ne0, ne1, s1, s2, n_dims, pos, freq_scale, ext_factor, - attn_factor, corr_dims, theta_scale, freq_factors); + x, dst, ne0, ne1, s1, s2, n_dims, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale, + freq_factors, row_indices, set_rows_stride); } else { rope_norm<<>>( - x, dst, ne0, ne1, s1, s2, n_dims, pos, freq_scale, ext_factor, - attn_factor, corr_dims, theta_scale, freq_factors); + x, dst, ne0, ne1, s1, s2, n_dims, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale, + freq_factors, row_indices, set_rows_stride); } } -template -static void rope_neox_cuda( - const T * x, T * dst, const int ne0, const int ne1, const int s1, const int s2, const int n_dims, const int nr, - const int32_t * pos, const float freq_scale, const float freq_base, const float ext_factor, const float attn_factor, - const rope_corr_dims corr_dims, const float * freq_factors, cudaStream_t stream) { +template +static void rope_neox_cuda(const T * x, + D * dst, + const int ne0, + const int ne1, + const int s1, + const int s2, + const int n_dims, + const int nr, + const int32_t * pos, + const float freq_scale, + const float freq_base, + const float ext_factor, + const float attn_factor, + const rope_corr_dims corr_dims, + const float * freq_factors, + const int64_t * row_indices, + const int set_rows_stride, + cudaStream_t stream) { GGML_ASSERT(ne0 % 2 == 0); const dim3 block_dims(1, CUDA_ROPE_BLOCK_SIZE, 1); const int n_blocks_x = (ne0 + 2*CUDA_ROPE_BLOCK_SIZE - 1) / (2*CUDA_ROPE_BLOCK_SIZE); @@ -274,13 +349,13 @@ static void rope_neox_cuda( const float theta_scale = powf(freq_base, -2.0f/n_dims); if (freq_factors == nullptr) { - rope_neox<<>>( - x, dst, ne0, ne1, s1, s2, n_dims, pos, freq_scale, ext_factor, - attn_factor, corr_dims, theta_scale, freq_factors); + rope_neox<<>>( + x, dst, ne0, ne1, s1, s2, n_dims, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale, + freq_factors, row_indices, set_rows_stride); } else { - rope_neox<<>>( - x, dst, ne0, ne1, s1, s2, n_dims, pos, freq_scale, ext_factor, - attn_factor, corr_dims, theta_scale, freq_factors); + rope_neox<<>>( + x, dst, ne0, ne1, s1, s2, n_dims, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale, + freq_factors, row_indices, set_rows_stride); } } @@ -333,7 +408,9 @@ static void rope_vision_cuda( } template -void ggml_cuda_op_rope_impl(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { +void ggml_cuda_op_rope_impl(ggml_backend_cuda_context & ctx, + ggml_tensor * dst, + const ggml_tensor * set_rows = nullptr) { const ggml_tensor * src0 = dst->src[0]; const ggml_tensor * src1 = dst->src[1]; const ggml_tensor * src2 = dst->src[2]; @@ -341,12 +418,25 @@ void ggml_cuda_op_rope_impl(ggml_backend_cuda_context & ctx, ggml_tensor * dst) const float * src0_d = (const float *)src0->data; const float * src1_d = (const float *)src1->data; - float * dst_d = (float *)dst->data; + void * dst_d = dst->data; + const int64_t * row_indices = nullptr; + ggml_type dst_type = dst->type; + int set_rows_stride = 0; + + if (set_rows != nullptr) { + GGML_ASSERT(forward); + dst_d = set_rows->data; + row_indices = (const int64_t *) set_rows->src[1]->data; + dst_type = set_rows->type; + set_rows_stride = set_rows->nb[1] / ggml_type_size(set_rows->type); + } cudaStream_t stream = ctx.stream(); GGML_ASSERT(src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16); GGML_ASSERT( dst->type == GGML_TYPE_F32 || dst->type == GGML_TYPE_F16); - GGML_ASSERT(src0->type == dst->type); + // When not fused, src0 and dst types must match + // When fused (ROPE+VIEW+SET_ROWS), src0 may be F32 and dst may be F16 + GGML_ASSERT(src0->type == dst->type || (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F16)); const int64_t ne00 = src0->ne[0]; // head dims const int64_t ne01 = src0->ne[1]; // num heads @@ -404,14 +494,18 @@ void ggml_cuda_op_rope_impl(ggml_backend_cuda_context & ctx, ggml_tensor * dst) // compute if (is_neox) { - if (src0->type == GGML_TYPE_F32) { - rope_neox_cuda( - (const float *) src0_d, (float *) dst_d, ne00, ne01, s01, s02, n_dims, nr, pos, freq_scale, - freq_base, ext_factor, attn_factor, corr_dims, freq_factors, stream); - } else if (src0->type == GGML_TYPE_F16) { - rope_neox_cuda( - (const half *) src0_d, (half *) dst_d, ne00, ne01, s01, s02, n_dims, nr, pos, freq_scale, - freq_base, ext_factor, attn_factor, corr_dims, freq_factors, stream); + if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F32) { + rope_neox_cuda((const float *) src0_d, (float *) dst_d, ne00, ne01, s01, s02, n_dims, + nr, pos, freq_scale, freq_base, ext_factor, attn_factor, corr_dims, + freq_factors, row_indices, set_rows_stride, stream); + } else if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F16) { + rope_neox_cuda((const float *) src0_d, (half *) dst_d, ne00, ne01, s01, s02, n_dims, + nr, pos, freq_scale, freq_base, ext_factor, attn_factor, corr_dims, + freq_factors, row_indices, set_rows_stride, stream); + } else if (src0->type == GGML_TYPE_F16 && dst_type == GGML_TYPE_F16) { + rope_neox_cuda((const half *) src0_d, (half *) dst_d, ne00, ne01, s01, s02, n_dims, nr, + pos, freq_scale, freq_base, ext_factor, attn_factor, corr_dims, + freq_factors, row_indices, set_rows_stride, stream); } else { GGML_ABORT("fatal error"); } @@ -440,14 +534,18 @@ void ggml_cuda_op_rope_impl(ggml_backend_cuda_context & ctx, ggml_tensor * dst) GGML_ABORT("fatal error"); } } else { - if (src0->type == GGML_TYPE_F32) { - rope_norm_cuda( - (const float *) src0_d, (float *) dst_d, ne00, ne01, s01, s02, n_dims, nr, pos, freq_scale, - freq_base, ext_factor, attn_factor, corr_dims, freq_factors, stream); - } else if (src0->type == GGML_TYPE_F16) { - rope_norm_cuda( - (const half *) src0_d, (half *) dst_d, ne00, ne01, s01, s02, n_dims, nr, pos, freq_scale, - freq_base, ext_factor, attn_factor, corr_dims, freq_factors, stream); + if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F32) { + rope_norm_cuda((const float *) src0_d, (float *) dst_d, ne00, ne01, s01, s02, n_dims, + nr, pos, freq_scale, freq_base, ext_factor, attn_factor, corr_dims, + freq_factors, row_indices, set_rows_stride, stream); + } else if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F16) { + rope_norm_cuda((const float *) src0_d, (half *) dst_d, ne00, ne01, s01, s02, n_dims, + nr, pos, freq_scale, freq_base, ext_factor, attn_factor, corr_dims, + freq_factors, row_indices, set_rows_stride, stream); + } else if (src0->type == GGML_TYPE_F16 && dst_type == GGML_TYPE_F16) { + rope_norm_cuda((const half *) src0_d, (half *) dst_d, ne00, ne01, s01, s02, n_dims, nr, + pos, freq_scale, freq_base, ext_factor, attn_factor, corr_dims, + freq_factors, row_indices, set_rows_stride, stream); } else { GGML_ABORT("fatal error"); } @@ -461,3 +559,7 @@ void ggml_cuda_op_rope(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { void ggml_cuda_op_rope_back(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { ggml_cuda_op_rope_impl(ctx, dst); } + +void ggml_cuda_op_rope_fused(ggml_backend_cuda_context & ctx, ggml_tensor * rope, ggml_tensor * set_rows) { + ggml_cuda_op_rope_impl(ctx, rope, set_rows); +} diff --git a/ggml/src/ggml-cuda/rope.cuh b/ggml/src/ggml-cuda/rope.cuh index 9139f3b22..72af086cd 100644 --- a/ggml/src/ggml-cuda/rope.cuh +++ b/ggml/src/ggml-cuda/rope.cuh @@ -5,3 +5,5 @@ void ggml_cuda_op_rope(ggml_backend_cuda_context & ctx, ggml_tensor * dst); void ggml_cuda_op_rope_back(ggml_backend_cuda_context & ctx, ggml_tensor * dst); + +void ggml_cuda_op_rope_fused(ggml_backend_cuda_context & ctx, ggml_tensor * dst, ggml_tensor * set_rows); diff --git a/ggml/src/ggml-hexagon/ggml-hexagon.cpp b/ggml/src/ggml-hexagon/ggml-hexagon.cpp index 7064b7486..cabd301ad 100644 --- a/ggml/src/ggml-hexagon/ggml-hexagon.cpp +++ b/ggml/src/ggml-hexagon/ggml-hexagon.cpp @@ -3156,26 +3156,17 @@ static inline bool op_reuse_src1(const ggml_tensor * op1, const ggml_tensor * op return (op0 && op0->src[1] == op1->src[1]); } +static inline bool is_compute_op(ggml_tensor *node) +{ + return !(ggml_op_is_empty(node->op) || ggml_is_empty(node)); +} + // scan the graph and figure out last compute op index static inline int last_compute_op(ggml_cgraph * graph) { - int last; + int last = 0; for (int i = 0; i < graph->n_nodes; ++i) { - ggml_tensor * node = graph->nodes[i]; - - switch (node->op) { - case GGML_OP_MUL_MAT: - case GGML_OP_MUL_MAT_ID: - case GGML_OP_MUL: - case GGML_OP_ADD: - case GGML_OP_SUB: - case GGML_OP_RMS_NORM: - case GGML_OP_GLU: - case GGML_OP_ADD_ID: - last = i; - break; - - default: - break; + if (is_compute_op(graph->nodes[i])) { + last = i; } } @@ -3194,6 +3185,10 @@ static ggml_status ggml_backend_hexagon_graph_compute(ggml_backend_t backend, gg for (int i = 0; i < graph->n_nodes; ++i) { ggml_tensor * node = graph->nodes[i]; + if (!is_compute_op(node)) { + continue; + } + uint32_t flags = 0; // skip quantizer if src1 is reused @@ -3245,14 +3240,6 @@ static ggml_status ggml_backend_hexagon_graph_compute(ggml_backend_t backend, gg ggml_hexagon_rope(node, flags); break; - // non-compute ops - case GGML_OP_NONE: - case GGML_OP_RESHAPE: - case GGML_OP_VIEW: - case GGML_OP_PERMUTE: - case GGML_OP_TRANSPOSE: - break; - default: GGML_ABORT("\nggml-hex: graph-compute %s is not supported\n", ggml_op_desc(node)); } diff --git a/ggml/src/ggml-hexagon/htp/binary-ops.c b/ggml/src/ggml-hexagon/htp/binary-ops.c index 92c0109d2..8ed7f67d9 100644 --- a/ggml/src/ggml-hexagon/htp/binary-ops.c +++ b/ggml/src/ggml-hexagon/htp/binary-ops.c @@ -34,6 +34,11 @@ static hvx_elemwise_f32_func func_table_HVX[] = { hvx_mul_f32, hvx_add_f32, static hvx_elemwise_f32_func func_table_HVX_opt[] = { hvx_mul_f32_opt, hvx_add_f32_opt, hvx_sub_f32_opt }; #define htp_binary_preamble \ + const struct htp_tensor * src0 = &octx->src0; \ + const struct htp_tensor * src1 = &octx->src1; \ + const struct htp_tensor * src2 = &octx->src2; \ + struct htp_tensor * dst = &octx->dst; \ + \ const uint32_t ne00 = src0->ne[0]; \ const uint32_t ne01 = src0->ne[1]; \ const uint32_t ne02 = src0->ne[2]; \ @@ -62,16 +67,15 @@ static hvx_elemwise_f32_func func_table_HVX_opt[] = { hvx_mul_f32_opt, hvx_add_f const uint32_t nb0 = dst->nb[0]; \ const uint32_t nb1 = dst->nb[1]; \ const uint32_t nb2 = dst->nb[2]; \ - const uint32_t nb3 = dst->nb[3]; + const uint32_t nb3 = dst->nb[3]; \ + \ + const uint32_t src0_nrows_per_thread = octx->src0_nrows_per_thread; -static void binary_job_f32_per_thread(const struct htp_tensor * src0, - const struct htp_tensor * src1, - struct htp_tensor * dst, - uint8_t * spad_data, - uint32_t nth, - uint32_t ith, - uint32_t src0_nrows_per_thread, - enum htp_op op) { +static void binary_job_f32_per_thread(struct htp_ops_context * octx, + uint8_t * spad_data, + uint32_t nth, + uint32_t ith, + enum htp_op op) { htp_binary_preamble; const size_t src0_row_size = nb01; @@ -107,16 +111,23 @@ static void binary_job_f32_per_thread(const struct htp_tensor * src0, uint8_t * restrict spad_data_th = spad_data + (ith * src0_row_size); - const uint32_t nr0 = ne00 / ne10; - const uint8_t * restrict src0_ptr = (const uint8_t *) src0->data + (src0_start_row * src0_row_size); uint8_t * restrict dst_ptr = (uint8_t *) dst->data + (src0_start_row * dst_row_size); const uint8_t * restrict data_src1 = (const uint8_t *) src1->data; - const uint8_t * restrict src1_ptr = NULL; + + const uint32_t ne02_ne01 = ne02 * ne01; for (uint32_t ir = src0_start_row; ir < src0_end_row; ir++) { - src1_ptr = data_src1 + (ir % src1_nrows) * src1_row_size; + const uint32_t i03 = fastdiv(ir, &octx->src0_div21); + const uint32_t i02 = fastdiv(ir - i03 * ne02_ne01, &octx->src0_div1); + const uint32_t i01 = (ir - i03 * ne02_ne01 - i02 * ne01); + + const uint32_t i13 = fastmodulo(i03, ne13, &octx->src1_div3); + const uint32_t i12 = fastmodulo(i02, ne12, &octx->src1_div2); + const uint32_t i11 = fastmodulo(i01, ne11, &octx->src1_div1); + + const uint8_t * restrict src1_ptr = data_src1 + i13 * nb13 + i12 * nb12 + i11 * src1_row_size; if (ir + 1 < src0_end_row) { htp_l2fetch(src0_ptr + ne00, 1, src0_row_size, src0_row_size); @@ -125,6 +136,7 @@ static void binary_job_f32_per_thread(const struct htp_tensor * src0, } } + const uint32_t nr0 = ne00 / ne10; if (nr0 > 1) { if ((1 == is_aligned) && (nr0 == ne00)) { hvx_bcast_fp32_a(spad_data_th, *(float *) src1_ptr, nr0); @@ -149,22 +161,17 @@ static void binary_job_f32_per_thread(const struct htp_tensor * src0, (unsigned) HAP_perf_qtimer_count_to_us(t2 - t1)); } -static void binary_add_id_job_f32_per_thread(const struct htp_tensor * src0, - const struct htp_tensor * src1, - const struct htp_tensor * src2, - struct htp_tensor * dst, - uint8_t * spad_data, - uint32_t nth, - uint32_t ith, - uint32_t src0_nrows_per_thread, - hvx_elemwise_f32_func func_HVX) { +static void binary_add_id_job_f32_per_thread(struct htp_ops_context * octx, + uint8_t * spad_data, + uint32_t nth, + uint32_t ith, + hvx_elemwise_f32_func func_HVX) { htp_binary_preamble; const size_t src0_row_size = nb01; const size_t src1_row_size = nb11; const size_t dst_row_size = nb1; - const uint32_t ne02_ne01 = ne02 * ne01; const uint32_t src0_nrows = ne01 * ne02 * ne03; // src0 rows const uint32_t src0_start_row = src0_nrows_per_thread * ith; @@ -187,10 +194,11 @@ static void binary_add_id_job_f32_per_thread(const struct htp_tensor * src0, const uint8_t * restrict data_src1 = (const uint8_t *) src1->data; uint8_t * restrict data_dst = (uint8_t *) dst->data; + const uint32_t ne02_ne01 = ne02 * ne01; for (uint32_t ir = src0_start_row; ir < src0_end_row; ir++) { // src0 indices - const uint32_t i03 = ir / ne02_ne01; - const uint32_t i02 = (ir - i03 * ne02_ne01) / ne01; + const uint32_t i03 = fastdiv(ir, &octx->src0_div21); + const uint32_t i02 = fastdiv(ir - i03 * ne02_ne01, &octx->src0_div1); const uint32_t i01 = (ir - i03 * ne02_ne01 - i02 * ne01); // src1 indices @@ -234,13 +242,11 @@ static void binary_job_dispatcher_f32(unsigned int n, unsigned int i, void * dat case HTP_OP_MUL: case HTP_OP_ADD: case HTP_OP_SUB: - binary_job_f32_per_thread(&octx->src0, &octx->src1, &octx->dst, octx->src1_spad.data, n, i, - octx->src0_nrows_per_thread, octx->op); + binary_job_f32_per_thread(octx, octx->src1_spad.data, n, i, octx->op); break; case HTP_OP_ADD_ID: - binary_add_id_job_f32_per_thread(&octx->src0, &octx->src1, &octx->src2, &octx->dst, octx->src0_spad.data, n, - i, octx->src0_nrows_per_thread, hvx_add_f32); + binary_add_id_job_f32_per_thread(octx, octx->src0_spad.data, n, i, hvx_add_f32); break; default: @@ -321,6 +327,16 @@ static int execute_op_binary_f32(struct htp_ops_context * octx) { octx->src0_nrows_per_thread = (src0_nrows + n_jobs - 1) / n_jobs; + octx->src0_div21 = init_fastdiv_values(src0->ne[2] * src0->ne[1]); + octx->src0_div3 = init_fastdiv_values(src0->ne[3]); + octx->src0_div2 = init_fastdiv_values(src0->ne[2]); + octx->src0_div1 = init_fastdiv_values(src0->ne[1]); + + octx->src1_div21 = init_fastdiv_values(src1->ne[2] * src1->ne[1]); + octx->src1_div3 = init_fastdiv_values(src1->ne[3]); + octx->src1_div2 = init_fastdiv_values(src1->ne[2]); + octx->src1_div1 = init_fastdiv_values(src1->ne[1]); + worker_pool_run_func(octx->ctx->worker_pool, binary_op_func, octx, n_jobs); } diff --git a/ggml/src/ggml-hexagon/htp/htp-msg.h b/ggml/src/ggml-hexagon/htp/htp-msg.h index f23d57880..9278f41f4 100644 --- a/ggml/src/ggml-hexagon/htp/htp-msg.h +++ b/ggml/src/ggml-hexagon/htp/htp-msg.h @@ -119,10 +119,10 @@ static const char * htp_type_name(uint32_t t) { #define HTP_MAX_DIMS 4 struct htp_tensor { - uint32_t data; // Buffer offset in the messages, and data pointer on the NSP - uint32_t type; // Data type - uint32_t ne[HTP_MAX_DIMS]; // Number of elements - uint32_t nb[HTP_MAX_DIMS]; // Stride in bytes (see ggml.h ggml_tensor) + uint32_t data; // Buffer offset in the messages, and data pointer on the NSP + uint32_t type; // Data type + uint32_t ne[HTP_MAX_DIMS]; // Number of elements + uint32_t nb[HTP_MAX_DIMS]; // Stride in bytes (see ggml.h ggml_tensor) }; #define HTP_MAX_OP_PARAMS 64 diff --git a/ggml/src/ggml-hexagon/htp/htp-ops.h b/ggml/src/ggml-hexagon/htp/htp-ops.h index 457231967..e87657436 100644 --- a/ggml/src/ggml-hexagon/htp/htp-ops.h +++ b/ggml/src/ggml-hexagon/htp/htp-ops.h @@ -4,6 +4,7 @@ #include "htp-ctx.h" #include "htp-msg.h" #include "worker-pool.h" +#include "ops-utils.h" #include #include @@ -38,6 +39,16 @@ struct htp_ops_context { uint32_t src0_nrows_per_thread; uint32_t src1_nrows_per_thread; + struct fastdiv_values src0_div1; // fastdiv values for ne1 + struct fastdiv_values src0_div2; // fastdiv values for ne2 + struct fastdiv_values src0_div3; // fastdiv values for ne3 + struct fastdiv_values src0_div21; // fastdiv values for ne2 * ne1 + + struct fastdiv_values src1_div1; // fastdiv values for ne1 + struct fastdiv_values src1_div2; // fastdiv values for ne2 + struct fastdiv_values src1_div3; // fastdiv values for ne3 + struct fastdiv_values src1_div21; // fastdiv values for ne2 * ne1 + uint32_t flags; }; diff --git a/ggml/src/ggml-hexagon/htp/ops-utils.h b/ggml/src/ggml-hexagon/htp/ops-utils.h index 302f16252..af9c3305f 100644 --- a/ggml/src/ggml-hexagon/htp/ops-utils.h +++ b/ggml/src/ggml-hexagon/htp/ops-utils.h @@ -31,6 +31,39 @@ static inline uint32_t htp_round_up(uint32_t n, uint32_t m) { return m * ((n + m - 1) / m); } +// See https://gmplib.org/~tege/divcnst-pldi94.pdf figure 4.1. +// Precompute mp (m' in the paper) and L such that division +// can be computed using a multiply (high 32b of 64b result) +// and a shift: +// +// n/d = (mulhi(n, mp) + n) >> L; +struct fastdiv_values { + uint32_t mp; + uint32_t l; +}; + +static inline struct fastdiv_values init_fastdiv_values(uint32_t d) { + struct fastdiv_values result = { 0, 0 }; + // compute L = ceil(log2(d)); + while (result.l < 32 && ((uint32_t) 1 << result.l) < d) { + ++(result.l); + } + + result.mp = (uint32_t) (((uint64_t) 1 << 32) * (((uint64_t) 1 << result.l) - d) / d + 1); + return result; +} + +static inline uint32_t fastdiv(uint32_t n, const struct fastdiv_values * vals) { + // Compute high 32 bits of n * mp + const uint32_t hi = (uint32_t) (((uint64_t) n * vals->mp) >> 32); // mulhi(n, mp) + // add n, apply bit shift + return (hi + n) >> vals->l; +} + +static inline uint32_t fastmodulo(uint32_t n, uint32_t d, const struct fastdiv_values * vals) { + return n - fastdiv(n, vals) * d; +} + static inline void htp_l2fetch(const void * p, uint32_t height, uint32_t width, uint32_t stride) { const uint64_t control = Q6_P_combine_RR(stride, Q6_R_combine_RlRl(width, height)); asm volatile(" l2fetch(%0,%1) " : : "r"(p), "r"(control)); diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 83d61fb4a..5d3eb9ccb 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -12700,6 +12700,12 @@ static bool ggml_vk_can_fuse_rms_norm_mul_rope(ggml_backend_vk_context * ctx, co return false; } + // conditions for pipeline creation + if (!(ctx->device->float_controls_rte_fp16 && + sizeof(vk_op_rms_norm_mul_rope_push_constants) <= ctx->device->properties.limits.maxPushConstantsSize)) { + return false; + } + return true; } diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp index b199e9462..650e40ec6 100644 --- a/src/llama-graph.cpp +++ b/src/llama-graph.cpp @@ -1592,9 +1592,10 @@ ggml_tensor * llm_graph_context::build_attn( int il) const { // these nodes are added to the graph together so that they are not reordered // by doing so, the number of splits in the graph is reduced + // expand k later to enable rope fusion which directly writes into k-v cache ggml_build_forward_expand(gf, q_cur); - ggml_build_forward_expand(gf, k_cur); ggml_build_forward_expand(gf, v_cur); + ggml_build_forward_expand(gf, k_cur); const auto * mctx_cur = inp->mctx; diff --git a/src/llama-vocab.cpp b/src/llama-vocab.cpp index 0c4e6901e..6a5eb0051 100644 --- a/src/llama-vocab.cpp +++ b/src/llama-vocab.cpp @@ -1238,7 +1238,7 @@ private: } private: uint32_t get_node(size_t index) { - if (index > xcda_array_size) { + if (index >= xcda_array_size) { throw std::runtime_error("Index out of array bounds in XCDA array!"); } return xcda_array[index]; diff --git a/tools/server/server.cpp b/tools/server/server.cpp index 6bd4be3cc..0b3c77879 100644 --- a/tools/server/server.cpp +++ b/tools/server/server.cpp @@ -684,7 +684,7 @@ struct server_task_result { } virtual bool is_stop() { // only used by server_task_result_cmpl_* - return false; + return true; } virtual int get_index() { return -1; @@ -3238,105 +3238,6 @@ struct server_context { queue_results.send(std::move(res)); } - // - // Functions to create new task(s) and receive result(s) - // - - void cancel_tasks(const std::unordered_set & id_tasks) { - std::vector cancel_tasks; - cancel_tasks.reserve(id_tasks.size()); - for (const auto & id_task : id_tasks) { - SRV_WRN("cancel task, id_task = %d\n", id_task); - - server_task task(SERVER_TASK_TYPE_CANCEL); - task.id_target = id_task; - queue_results.remove_waiting_task_id(id_task); - cancel_tasks.push_back(std::move(task)); - } - // push to beginning of the queue, so it has highest priority - queue_tasks.post(std::move(cancel_tasks), true); - } - - // receive the results from task(s) - void receive_multi_results( - const std::unordered_set & id_tasks, - const std::function&)> & result_handler, - const std::function & error_handler, - const std::function & is_connection_closed) { - std::vector results(id_tasks.size()); - for (int i = 0; i < (int)id_tasks.size(); i++) { - server_task_result_ptr result = queue_results.recv_with_timeout(id_tasks, HTTP_POLLING_SECONDS); - - if (is_connection_closed()) { - cancel_tasks(id_tasks); - return; - } - - if (result == nullptr) { - i--; // retry - continue; - } - - if (result->is_error()) { - error_handler(result->to_json()); - cancel_tasks(id_tasks); - return; - } - - GGML_ASSERT( - dynamic_cast(result.get()) != nullptr - || dynamic_cast(result.get()) != nullptr - || dynamic_cast(result.get()) != nullptr - ); - const size_t idx = result->get_index(); - GGML_ASSERT(idx < results.size() && "index out of range"); - results[idx] = std::move(result); - } - result_handler(results); - } - - // receive the results from task(s), in stream mode - void receive_cmpl_results_stream( - const std::unordered_set & id_tasks, - const std::function & result_handler, - const std::function & error_handler, - const std::function & is_connection_closed) { - size_t n_finished = 0; - while (true) { - server_task_result_ptr result = queue_results.recv_with_timeout(id_tasks, HTTP_POLLING_SECONDS); - - if (is_connection_closed()) { - cancel_tasks(id_tasks); - return; - } - - if (result == nullptr) { - continue; // retry - } - - if (result->is_error()) { - error_handler(result->to_json()); - cancel_tasks(id_tasks); - return; - } - - GGML_ASSERT( - dynamic_cast(result.get()) != nullptr - || dynamic_cast(result.get()) != nullptr - ); - if (!result_handler(result)) { - cancel_tasks(id_tasks); - break; - } - - if (result->is_stop()) { - if (++n_finished == id_tasks.size()) { - break; - } - } - } - } - // // Functions to process the task // @@ -4418,6 +4319,104 @@ struct server_context { } }; +// generator-like API for server responses, support pooling connection state and aggregating results +struct server_response_reader { + std::unordered_set id_tasks; + server_context & ctx_server; + size_t received_count = 0; + bool cancelled = false; + + server_response_reader(server_context & ctx_server) : ctx_server(ctx_server) {} + ~server_response_reader() { + stop(); + } + + void post_tasks(std::vector && tasks) { + id_tasks = server_task::get_list_id(tasks); + ctx_server.queue_results.add_waiting_tasks(tasks); + ctx_server.queue_tasks.post(std::move(tasks)); + } + + bool has_next() { + return !cancelled && received_count < id_tasks.size(); + } + + // return nullptr if should_stop() is true before receiving a result + // note: if one error is received, it will stop further processing and return error result + server_task_result_ptr next(const std::function & should_stop) { + while (true) { + server_task_result_ptr result = ctx_server.queue_results.recv_with_timeout(id_tasks, HTTP_POLLING_SECONDS); + if (result == nullptr) { + // timeout, check stop condition + if (should_stop()) { + SRV_DBG("%s", "stopping wait for next result due to should_stop condition\n"); + return nullptr; + } + } else { + if (result->is_error()) { + stop(); // cancel remaining tasks + SRV_DBG("%s", "received error result, stopping further processing\n"); + return result; + } + if (result->is_stop()) { + received_count++; + } + return result; + } + } + + // should not reach here + } + + struct batch_response { + bool is_terminated = false; // if true, indicates that processing was stopped before all results were received + std::vector results; + server_task_result_ptr error; // nullptr if no error + }; + + batch_response wait_for_all(const std::function & should_stop) { + batch_response batch_res; + batch_res.results.resize(id_tasks.size()); + while (has_next()) { + auto res = next(should_stop); + if (res == nullptr) { + batch_res.is_terminated = true; + return batch_res; + } + if (res->is_error()) { + batch_res.error = std::move(res); + return batch_res; + } + const size_t idx = res->get_index(); + GGML_ASSERT(idx < batch_res.results.size() && "index out of range"); + GGML_ASSERT(batch_res.results[idx] == nullptr && "duplicate result received"); + batch_res.results[idx] = std::move(res); + } + return batch_res; + } + + void stop() { + ctx_server.queue_results.remove_waiting_task_ids(id_tasks); + if (has_next() && !cancelled) { + // if tasks is not finished yet, cancel them + cancelled = true; + std::vector cancel_tasks; + cancel_tasks.reserve(id_tasks.size()); + for (const auto & id_task : id_tasks) { + SRV_WRN("cancel task, id_task = %d\n", id_task); + server_task task(SERVER_TASK_TYPE_CANCEL); + task.id_target = id_task; + ctx_server.queue_results.remove_waiting_task_id(id_task); + cancel_tasks.push_back(std::move(task)); + } + // push to beginning of the queue, so it has highest priority + ctx_server.queue_tasks.post(std::move(cancel_tasks), true); + } else { + SRV_DBG("%s", "all tasks already finished, no need to cancel\n"); + } + } +}; + static void log_server_request(const httplib::Request & req, const httplib::Response & res) { // skip GH copilot requests when using default port if (req.path == "/v1/health") { @@ -4432,6 +4431,17 @@ static void log_server_request(const httplib::Request & req, const httplib::Resp SRV_DBG("response: %s\n", res.body.c_str()); } +static void res_error(httplib::Response & res, const json & error_data) { + json final_response {{"error", error_data}}; + res.set_content(safe_json_to_str(final_response), MIMETYPE_JSON); + res.status = json_value(error_data, "code", 500); +} + +static void res_ok(httplib::Response & res, const json & data) { + res.set_content(safe_json_to_str(data), MIMETYPE_JSON); + res.status = 200; +} + std::function shutdown_handler; std::atomic_flag is_terminating = ATOMIC_FLAG_INIT; @@ -4501,19 +4511,7 @@ int main(int argc, char ** argv) { svr->set_default_headers({{"Server", "llama.cpp"}}); svr->set_logger(log_server_request); - - auto res_error = [](httplib::Response & res, const json & error_data) { - json final_response {{"error", error_data}}; - res.set_content(safe_json_to_str(final_response), MIMETYPE_JSON); - res.status = json_value(error_data, "code", 500); - }; - - auto res_ok = [](httplib::Response & res, const json & data) { - res.set_content(safe_json_to_str(data), MIMETYPE_JSON); - res.status = 200; - }; - - svr->set_exception_handler([&res_error](const httplib::Request &, httplib::Response & res, const std::exception_ptr & ep) { + svr->set_exception_handler([](const httplib::Request &, httplib::Response & res, const std::exception_ptr & ep) { std::string message; try { std::rethrow_exception(ep); @@ -4532,7 +4530,7 @@ int main(int argc, char ** argv) { } }); - svr->set_error_handler([&res_error](const httplib::Request &, httplib::Response & res) { + svr->set_error_handler([](const httplib::Request &, httplib::Response & res) { if (res.status == 404) { res_error(res, format_error_response("File Not Found", ERROR_TYPE_NOT_FOUND)); } @@ -4562,7 +4560,7 @@ int main(int argc, char ** argv) { // Middlewares // - auto middleware_validate_api_key = [¶ms, &res_error](const httplib::Request & req, httplib::Response & res) { + auto middleware_validate_api_key = [¶ms](const httplib::Request & req, httplib::Response & res) { static const std::unordered_set public_endpoints = { "/health", "/v1/health", @@ -4600,7 +4598,7 @@ int main(int argc, char ** argv) { return false; }; - auto middleware_server_state = [&res_error, &state](const httplib::Request & req, httplib::Response & res) { + auto middleware_server_state = [&state](const httplib::Request & req, httplib::Response & res) { server_state current_state = state.load(); if (current_state == SERVER_STATE_LOADING_MODEL) { auto tmp = string_split(req.path, '.'); @@ -4788,7 +4786,7 @@ int main(int argc, char ** argv) { res.status = 200; // HTTP OK }; - const auto handle_slots_save = [&ctx_server, &res_error, &res_ok, ¶ms](const httplib::Request & req, httplib::Response & res, int id_slot) { + const auto handle_slots_save = [&ctx_server, ¶ms](const httplib::Request & req, httplib::Response & res, int id_slot) { json request_data = json::parse(req.body); std::string filename = request_data.at("filename"); if (!fs_validate_filename(filename)) { @@ -4820,7 +4818,7 @@ int main(int argc, char ** argv) { res_ok(res, result->to_json()); }; - const auto handle_slots_restore = [&ctx_server, &res_error, &res_ok, ¶ms](const httplib::Request & req, httplib::Response & res, int id_slot) { + const auto handle_slots_restore = [&ctx_server, ¶ms](const httplib::Request & req, httplib::Response & res, int id_slot) { json request_data = json::parse(req.body); std::string filename = request_data.at("filename"); if (!fs_validate_filename(filename)) { @@ -4853,7 +4851,7 @@ int main(int argc, char ** argv) { res_ok(res, result->to_json()); }; - const auto handle_slots_erase = [&ctx_server, &res_error, &res_ok](const httplib::Request & /* req */, httplib::Response & res, int id_slot) { + const auto handle_slots_erase = [&ctx_server](const httplib::Request & /* req */, httplib::Response & res, int id_slot) { int task_id = ctx_server.queue_tasks.get_new_id(); { server_task task(SERVER_TASK_TYPE_SLOT_ERASE); @@ -4876,7 +4874,7 @@ int main(int argc, char ** argv) { res_ok(res, result->to_json()); }; - const auto handle_slots_action = [¶ms, &res_error, &handle_slots_save, &handle_slots_restore, &handle_slots_erase](const httplib::Request & req, httplib::Response & res) { + const auto handle_slots_action = [¶ms, &handle_slots_save, &handle_slots_restore, &handle_slots_erase](const httplib::Request & req, httplib::Response & res) { if (params.slot_save_path.empty()) { res_error(res, format_error_response("This server does not support slots action. Start it with `--slot-save-path`", ERROR_TYPE_NOT_SUPPORTED)); return; @@ -4905,7 +4903,7 @@ int main(int argc, char ** argv) { } }; - const auto handle_props = [¶ms, &ctx_server, &res_ok](const httplib::Request &, httplib::Response & res) { + const auto handle_props = [¶ms, &ctx_server](const httplib::Request &, httplib::Response & res) { json default_generation_settings_for_props; { @@ -4947,7 +4945,7 @@ int main(int argc, char ** argv) { res_ok(res, data); }; - const auto handle_props_change = [&ctx_server, &res_error, &res_ok](const httplib::Request & req, httplib::Response & res) { + const auto handle_props_change = [&ctx_server](const httplib::Request & req, httplib::Response & res) { if (!ctx_server.params_base.endpoint_props) { res_error(res, format_error_response("This server does not support changing global properties. Start it with `--props`", ERROR_TYPE_NOT_SUPPORTED)); return; @@ -4960,7 +4958,7 @@ int main(int argc, char ** argv) { res_ok(res, {{ "success", true }}); }; - const auto handle_api_show = [&ctx_server, &res_ok](const httplib::Request &, httplib::Response & res) { + const auto handle_api_show = [&ctx_server](const httplib::Request &, httplib::Response & res) { bool has_mtmd = ctx_server.mctx != nullptr; json data = { { @@ -4991,7 +4989,7 @@ int main(int argc, char ** argv) { // handle completion-like requests (completion, chat, infill) // we can optionally provide a custom format for partial results and final results - const auto handle_completions_impl = [&ctx_server, &res_error, &res_ok]( + const auto handle_completions_impl = [&ctx_server]( server_task_type type, json & data, const std::vector & files, @@ -5001,7 +4999,10 @@ int main(int argc, char ** argv) { GGML_ASSERT(type == SERVER_TASK_TYPE_COMPLETION || type == SERVER_TASK_TYPE_INFILL); auto completion_id = gen_chatcmplid(); - std::unordered_set task_ids; + // need to store the reader as a pointer, so that it won't be destroyed when the handle returns + // use shared_ptr as it's shared between the chunked_content_provider() and on_complete() + const auto rd = std::make_shared(ctx_server); + try { std::vector tasks; @@ -5019,17 +5020,8 @@ int main(int argc, char ** argv) { // Everything else, including multimodal completions. inputs = tokenize_input_prompts(ctx_server.vocab, ctx_server.mctx, prompt, true, true); } - const size_t n_ctx_slot = ctx_server.slots.front().n_ctx; tasks.reserve(inputs.size()); for (size_t i = 0; i < inputs.size(); i++) { - auto n_prompt_tokens = inputs[i].size(); - if (n_prompt_tokens >= n_ctx_slot) { - json error_data = format_error_response("the request exceeds the available context size, try increasing it", ERROR_TYPE_EXCEED_CONTEXT_SIZE); - error_data["n_prompt_tokens"] = n_prompt_tokens; - error_data["n_ctx"] = n_ctx_slot; - res_error(res, error_data); - return; - } server_task task = server_task(type); task.id = ctx_server.queue_tasks.get_new_id(); @@ -5050,9 +5042,7 @@ int main(int argc, char ** argv) { tasks.push_back(std::move(task)); } - task_ids = server_task::get_list_id(tasks); - ctx_server.queue_results.add_waiting_tasks(tasks); - ctx_server.queue_tasks.post(std::move(tasks)); + rd->post_tasks(std::move(tasks)); } catch (const std::exception & e) { res_error(res, format_error_response(e.what(), ERROR_TYPE_INVALID_REQUEST)); return; @@ -5061,54 +5051,95 @@ int main(int argc, char ** argv) { bool stream = json_value(data, "stream", false); if (!stream) { - ctx_server.receive_multi_results(task_ids, [&](std::vector & results) { - if (results.size() == 1) { - // single result - res_ok(res, results[0]->to_json()); - } else { - // multiple results (multitask) - json arr = json::array(); - for (auto & res : results) { - arr.push_back(res->to_json()); - } - res_ok(res, arr); + // non-stream, wait for the results + auto all_results = rd->wait_for_all(is_connection_closed); + if (all_results.is_terminated) { + return; // connection is closed + } else if (all_results.error) { + res_error(res, all_results.error->to_json()); + return; + } else { + json arr = json::array(); + for (auto & res : all_results.results) { + GGML_ASSERT(dynamic_cast(res.get()) != nullptr); + arr.push_back(res->to_json()); } - }, [&](const json & error_data) { - res_error(res, error_data); - }, is_connection_closed); + // if single request, return single object instead of array + res_ok(res, arr.size() == 1 ? arr[0] : arr); + } - ctx_server.queue_results.remove_waiting_task_ids(task_ids); } else { - const auto chunked_content_provider = [task_ids, &ctx_server, oaicompat](size_t, httplib::DataSink & sink) { - ctx_server.receive_cmpl_results_stream(task_ids, [&](server_task_result_ptr & result) -> bool { - json res_json = result->to_json(); - if (res_json.is_array()) { - for (const auto & res : res_json) { - if (!server_sent_event(sink, res)) { - // sending failed (HTTP connection closed), cancel the generation - return false; - } - } - return true; - } else { - return server_sent_event(sink, res_json); + // in streaming mode, the first error must be treated as non-stream response + // this is to match the OAI API behavior + // ref: https://github.com/ggml-org/llama.cpp/pull/16486#discussion_r2419657309 + server_task_result_ptr first_result = rd->next(is_connection_closed); + if (first_result == nullptr) { + return; // connection is closed + } else if (first_result->is_error()) { + res_error(res, first_result->to_json()); + return; + } else { + GGML_ASSERT( + dynamic_cast(first_result.get()) != nullptr + || dynamic_cast(first_result.get()) != nullptr + ); + } + + // next responses are streamed + json first_result_json = first_result->to_json(); + const auto chunked_content_provider = [first_result_json, rd, oaicompat](size_t, httplib::DataSink & sink) mutable -> bool { + // flush the first result as it's not an error + if (!first_result_json.empty()) { + if (!server_sent_event(sink, first_result_json)) { + sink.done(); + return false; // sending failed, go to on_complete() } - }, [&](const json & error_data) { - server_sent_event(sink, json{{"error", error_data}}); - }, [&sink]() { - // note: do not use req.is_connection_closed here because req is already destroyed - return !sink.is_writable(); - }); - if (oaicompat != OAICOMPAT_TYPE_NONE) { - static const std::string ev_done = "data: [DONE]\n\n"; - sink.write(ev_done.data(), ev_done.size()); + first_result_json.clear(); // mark as sent } - sink.done(); - return false; + + // receive subsequent results + auto result = rd->next([&sink]{ return !sink.is_writable(); }); + if (result == nullptr) { + sink.done(); + return false; // connection is closed, go to on_complete() + } + + // send the results + json res_json = result->to_json(); + bool ok = false; + if (result->is_error()) { + ok = server_sent_event(sink, json {{ "error", result->to_json() }}); + sink.done(); + return false; // go to on_complete() + } else { + GGML_ASSERT( + dynamic_cast(result.get()) != nullptr + || dynamic_cast(result.get()) != nullptr + ); + ok = server_sent_event(sink, res_json); + } + + if (!ok) { + sink.done(); + return false; // sending failed, go to on_complete() + } + + // check if there is more data + if (!rd->has_next()) { + if (oaicompat != OAICOMPAT_TYPE_NONE) { + static const std::string ev_done = "data: [DONE]\n\n"; + sink.write(ev_done.data(), ev_done.size()); + } + sink.done(); + return false; // no more data, go to on_complete() + } + + // has next data, continue + return true; }; - auto on_complete = [task_ids, &ctx_server] (bool) { - ctx_server.queue_results.remove_waiting_task_ids(task_ids); + auto on_complete = [rd](bool) { + rd->stop(); }; res.set_chunked_content_provider("text/event-stream", chunked_content_provider, on_complete); @@ -5139,7 +5170,7 @@ int main(int argc, char ** argv) { OAICOMPAT_TYPE_COMPLETION); }; - const auto handle_infill = [&ctx_server, &res_error, &handle_completions_impl](const httplib::Request & req, httplib::Response & res) { + const auto handle_infill = [&ctx_server, &handle_completions_impl](const httplib::Request & req, httplib::Response & res) { // check model compatibility std::string err; if (llama_vocab_fim_pre(ctx_server.vocab) == LLAMA_TOKEN_NULL) { @@ -5238,7 +5269,7 @@ int main(int argc, char ** argv) { }; // same with handle_chat_completions, but without inference part - const auto handle_apply_template = [&ctx_server, &res_ok](const httplib::Request & req, httplib::Response & res) { + const auto handle_apply_template = [&ctx_server](const httplib::Request & req, httplib::Response & res) { auto body = json::parse(req.body); std::vector files; // dummy, unused json data = oaicompat_chat_params_parse( @@ -5248,7 +5279,7 @@ int main(int argc, char ** argv) { res_ok(res, {{ "prompt", std::move(data.at("prompt")) }}); }; - const auto handle_models = [¶ms, &ctx_server, &state, &res_ok](const httplib::Request &, httplib::Response & res) { + const auto handle_models = [¶ms, &ctx_server, &state](const httplib::Request &, httplib::Response & res) { server_state current_state = state.load(); json model_meta = nullptr; if (current_state == SERVER_STATE_READY) { @@ -5293,7 +5324,7 @@ int main(int argc, char ** argv) { res_ok(res, models); }; - const auto handle_tokenize = [&ctx_server, &res_ok](const httplib::Request & req, httplib::Response & res) { + const auto handle_tokenize = [&ctx_server](const httplib::Request & req, httplib::Response & res) { const json body = json::parse(req.body); json tokens_response = json::array(); @@ -5334,7 +5365,7 @@ int main(int argc, char ** argv) { res_ok(res, data); }; - const auto handle_detokenize = [&ctx_server, &res_ok](const httplib::Request & req, httplib::Response & res) { + const auto handle_detokenize = [&ctx_server](const httplib::Request & req, httplib::Response & res) { const json body = json::parse(req.body); std::string content; @@ -5347,7 +5378,7 @@ int main(int argc, char ** argv) { res_ok(res, data); }; - const auto handle_embeddings_impl = [&ctx_server, &res_error, &res_ok](const httplib::Request & req, httplib::Response & res, oaicompat_type oaicompat) { + const auto handle_embeddings_impl = [&ctx_server](const httplib::Request & req, httplib::Response & res, oaicompat_type oaicompat) { if (!ctx_server.params_base.embedding) { res_error(res, format_error_response("This server does not support embeddings. Start it with `--embeddings`", ERROR_TYPE_NOT_SUPPORTED)); return; @@ -5402,8 +5433,7 @@ int main(int argc, char ** argv) { // create and queue the task json responses = json::array(); - bool error = false; - std::unordered_set task_ids; + server_response_reader rd(ctx_server); { std::vector tasks; for (size_t i = 0; i < tokenized_prompts.size(); i++) { @@ -5419,27 +5449,23 @@ int main(int argc, char ** argv) { tasks.push_back(std::move(task)); } - - task_ids = server_task::get_list_id(tasks); - ctx_server.queue_results.add_waiting_tasks(tasks); - ctx_server.queue_tasks.post(std::move(tasks)); + rd.post_tasks(std::move(tasks)); } - // get the result - ctx_server.receive_multi_results(task_ids, [&](std::vector & results) { - for (auto & res : results) { + // wait for the results + auto all_results = rd.wait_for_all(req.is_connection_closed); + + // collect results + if (all_results.is_terminated) { + return; // connection is closed + } else if (all_results.error) { + res_error(res, all_results.error->to_json()); + return; + } else { + for (auto & res : all_results.results) { GGML_ASSERT(dynamic_cast(res.get()) != nullptr); responses.push_back(res->to_json()); } - }, [&](const json & error_data) { - res_error(res, error_data); - error = true; - }, req.is_connection_closed); - - ctx_server.queue_results.remove_waiting_task_ids(task_ids); - - if (error) { - return; } // write JSON response @@ -5457,7 +5483,7 @@ int main(int argc, char ** argv) { handle_embeddings_impl(req, res, OAICOMPAT_TYPE_EMBEDDING); }; - const auto handle_rerank = [&ctx_server, &res_error, &res_ok](const httplib::Request & req, httplib::Response & res) { + const auto handle_rerank = [&ctx_server](const httplib::Request & req, httplib::Response & res) { if (!ctx_server.params_base.embedding || ctx_server.params_base.pooling_type != LLAMA_POOLING_TYPE_RANK) { res_error(res, format_error_response("This server does not support reranking. Start it with `--reranking`", ERROR_TYPE_NOT_SUPPORTED)); return; @@ -5493,8 +5519,7 @@ int main(int argc, char ** argv) { // create and queue the task json responses = json::array(); - bool error = false; - std::unordered_set task_ids; + server_response_reader rd(ctx_server); { std::vector tasks; tasks.reserve(documents.size()); @@ -5506,24 +5531,23 @@ int main(int argc, char ** argv) { task.tokens = std::move(tmp); tasks.push_back(std::move(task)); } - - task_ids = server_task::get_list_id(tasks); - ctx_server.queue_results.add_waiting_tasks(tasks); - ctx_server.queue_tasks.post(std::move(tasks)); + rd.post_tasks(std::move(tasks)); } - ctx_server.receive_multi_results(task_ids, [&](std::vector & results) { - for (auto & res : results) { + // wait for the results + auto all_results = rd.wait_for_all(req.is_connection_closed); + + // collect results + if (all_results.is_terminated) { + return; // connection is closed + } else if (all_results.error) { + res_error(res, all_results.error->to_json()); + return; + } else { + for (auto & res : all_results.results) { GGML_ASSERT(dynamic_cast(res.get()) != nullptr); responses.push_back(res->to_json()); } - }, [&](const json & error_data) { - res_error(res, error_data); - error = true; - }, req.is_connection_closed); - - if (error) { - return; } // write JSON response diff --git a/tools/server/utils.hpp b/tools/server/utils.hpp index 2bce2f4a4..b1ecc5af5 100644 --- a/tools/server/utils.hpp +++ b/tools/server/utils.hpp @@ -9,14 +9,6 @@ #include "mtmd-helper.h" #include "chat.h" -// increase max payload length to allow use of larger context size -#define CPPHTTPLIB_FORM_URL_ENCODED_PAYLOAD_MAX_LENGTH 1048576 -// increase backlog size to avoid connection resets for >> 1 slots -#define CPPHTTPLIB_LISTEN_BACKLOG 512 -// increase max URI length to handle longer prompts in query string -#define CPPHTTPLIB_REQUEST_URI_MAX_LENGTH 32768 -// disable Nagle's algorithm -#define CPPHTTPLIB_TCP_NODELAY true #include #define JSON_ASSERT GGML_ASSERT @@ -461,15 +453,29 @@ static std::string tokens_to_output_formatted_string(const llama_context * ctx, return out; } +// note: if data is a json array, it will be sent as multiple events, one per item static bool server_sent_event(httplib::DataSink & sink, const json & data) { - const std::string str = - "data: " + - data.dump(-1, ' ', false, json::error_handler_t::replace) + - "\n\n"; // required by RFC 8895 - A message is terminated by a blank line (two line terminators in a row). + static auto send_single = [](httplib::DataSink & sink, const json & data) -> bool { + const std::string str = + "data: " + + data.dump(-1, ' ', false, json::error_handler_t::replace) + + "\n\n"; // required by RFC 8895 - A message is terminated by a blank line (two line terminators in a row). - LOG_DBG("data stream, to_send: %s", str.c_str()); + LOG_DBG("data stream, to_send: %s", str.c_str()); + return sink.write(str.c_str(), str.size()); + }; - return sink.write(str.c_str(), str.size()); + if (data.is_array()) { + for (const auto & item : data) { + if (!send_single(sink, item)) { + return false; + } + } + } else { + return send_single(sink, data); + } + + return true; } // diff --git a/tools/server/webui/.storybook/preview.ts b/tools/server/webui/.storybook/preview.ts index fb91386af..8d530e43e 100644 --- a/tools/server/webui/.storybook/preview.ts +++ b/tools/server/webui/.storybook/preview.ts @@ -11,8 +11,16 @@ const preview: Preview = { date: /Date$/i } }, + backgrounds: { disable: true + }, + + a11y: { + // 'todo' - show a11y violations in the test UI only + // 'error' - fail CI on a11y violations + // 'off' - skip a11y checks entirely + test: 'todo' } }, decorators: [ diff --git a/tools/server/webui/.storybook/vitest.setup.ts b/tools/server/webui/.storybook/vitest.setup.ts index e0c1753c8..147157289 100644 --- a/tools/server/webui/.storybook/vitest.setup.ts +++ b/tools/server/webui/.storybook/vitest.setup.ts @@ -1,8 +1,9 @@ +import * as a11yAddonAnnotations from '@storybook/addon-a11y/preview'; import { setProjectAnnotations } from '@storybook/sveltekit'; import * as previewAnnotations from './preview'; import { beforeAll } from 'vitest'; -const project = setProjectAnnotations([previewAnnotations]); +const project = setProjectAnnotations([a11yAddonAnnotations, previewAnnotations]); beforeAll(async () => { if (project.beforeAll) { diff --git a/tools/server/webui/package-lock.json b/tools/server/webui/package-lock.json index 8fab38f6f..a11b87ad5 100644 --- a/tools/server/webui/package-lock.json +++ b/tools/server/webui/package-lock.json @@ -22,20 +22,20 @@ "unist-util-visit": "^5.0.0" }, "devDependencies": { - "@chromatic-com/storybook": "^4.0.1", + "@chromatic-com/storybook": "^4.1.2", "@eslint/compat": "^1.2.5", "@eslint/js": "^9.18.0", "@internationalized/date": "^3.8.2", "@lucide/svelte": "^0.515.0", "@playwright/test": "^1.49.1", - "@storybook/addon-a11y": "^9.0.17", - "@storybook/addon-docs": "^9.0.17", - "@storybook/addon-svelte-csf": "^5.0.7", - "@storybook/addon-vitest": "^9.0.17", - "@storybook/sveltekit": "^9.0.17", - "@sveltejs/adapter-static": "^3.0.8", - "@sveltejs/kit": "^2.22.0", - "@sveltejs/vite-plugin-svelte": "^6.0.0", + "@storybook/addon-a11y": "^10.0.7", + "@storybook/addon-docs": "^10.0.7", + "@storybook/addon-svelte-csf": "^5.0.10", + "@storybook/addon-vitest": "^10.0.7", + "@storybook/sveltekit": "^10.0.7", + "@sveltejs/adapter-static": "^3.0.10", + "@sveltejs/kit": "^2.48.4", + "@sveltejs/vite-plugin-svelte": "^6.2.1", "@tailwindcss/forms": "^0.5.9", "@tailwindcss/typography": "^0.5.15", "@tailwindcss/vite": "^4.0.0", @@ -46,21 +46,21 @@ "dexie": "^4.0.11", "eslint": "^9.18.0", "eslint-config-prettier": "^10.0.1", - "eslint-plugin-storybook": "^9.0.17", + "eslint-plugin-storybook": "^10.0.7", "eslint-plugin-svelte": "^3.0.0", "fflate": "^0.8.2", "globals": "^16.0.0", "http-server": "^14.1.1", "mdast": "^3.0.0", "mdsvex": "^0.12.3", - "playwright": "^1.53.0", + "playwright": "^1.56.1", "prettier": "^3.4.2", "prettier-plugin-svelte": "^3.3.3", "prettier-plugin-tailwindcss": "^0.6.11", "rehype-katex": "^7.0.1", "remark-math": "^6.0.0", "sass": "^1.93.3", - "storybook": "^9.0.17", + "storybook": "^10.0.7", "svelte": "^5.0.0", "svelte-check": "^4.0.0", "tailwind-merge": "^3.3.1", @@ -71,7 +71,7 @@ "typescript-eslint": "^8.20.0", "unified": "^11.0.5", "uuid": "^13.0.0", - "vite": "^7.0.4", + "vite": "^7.2.2", "vite-plugin-devtools-json": "^0.2.0", "vitest": "^3.2.3", "vitest-browser-svelte": "^0.1.0" @@ -133,9 +133,9 @@ } }, "node_modules/@chromatic-com/storybook": { - "version": "4.0.1", - "resolved": "https://registry.npmjs.org/@chromatic-com/storybook/-/storybook-4.0.1.tgz", - "integrity": "sha512-GQXe5lyZl3yLewLJQyFXEpOp2h+mfN2bPrzYaOFNCJjO4Js9deKbRHTOSaiP2FRwZqDLdQwy2+SEGeXPZ94yYw==", + "version": "4.1.2", + "resolved": "https://registry.npmjs.org/@chromatic-com/storybook/-/storybook-4.1.2.tgz", + "integrity": "sha512-QAWGtHwib0qsP5CcO64aJCF75zpFgpKK3jNpxILzQiPK3sVo4EmnVGJVdwcZWpWrGdH8E4YkncGoitw4EXzKMg==", "dev": true, "license": "MIT", "dependencies": { @@ -150,7 +150,7 @@ "yarn": ">=1.22.18" }, "peerDependencies": { - "storybook": "^0.0.0-0 || ^9.0.0 || ^9.1.0-0" + "storybook": "^0.0.0-0 || ^9.0.0 || ^9.1.0-0 || ^9.2.0-0 || ^10.0.0-0 || ^10.1.0-0 || ^10.2.0-0 || ^10.3.0-0" } }, "node_modules/@esbuild/aix-ppc64": { @@ -894,6 +894,17 @@ "@jridgewell/trace-mapping": "^0.3.24" } }, + "node_modules/@jridgewell/remapping": { + "version": "2.3.5", + "resolved": "https://registry.npmjs.org/@jridgewell/remapping/-/remapping-2.3.5.tgz", + "integrity": "sha512-LI9u/+laYG4Ds1TDKSJW2YPrIlcVYOwi2fUC6xB43lueCjgxV4lffOCZCtYFiH6TNOX+tQKXx97T4IKHbhyHEQ==", + "dev": true, + "license": "MIT", + "dependencies": { + "@jridgewell/gen-mapping": "^0.3.5", + "@jridgewell/trace-mapping": "^0.3.24" + } + }, "node_modules/@jridgewell/resolve-uri": { "version": "3.1.2", "resolved": "https://registry.npmjs.org/@jridgewell/resolve-uri/-/resolve-uri-3.1.2.tgz", @@ -1502,13 +1513,13 @@ } }, "node_modules/@playwright/test": { - "version": "1.54.1", - "resolved": "https://registry.npmjs.org/@playwright/test/-/test-1.54.1.tgz", - "integrity": "sha512-FS8hQ12acieG2dYSksmLOF7BNxnVf2afRJdCuM1eMSxj6QTSE6G4InGF7oApGgDb65MX7AwMVlIkpru0yZA4Xw==", + "version": "1.56.1", + "resolved": "https://registry.npmjs.org/@playwright/test/-/test-1.56.1.tgz", + "integrity": "sha512-vSMYtL/zOcFpvJCW71Q/OEGQb7KYBPAdKh35WNSkaZA75JlAO8ED8UN6GUNTm3drWomcbcqRPFqQbLae8yBTdg==", "dev": true, "license": "Apache-2.0", "dependencies": { - "playwright": "1.54.1" + "playwright": "1.56.1" }, "bin": { "playwright": "cli.js" @@ -1812,9 +1823,9 @@ "license": "MIT" }, "node_modules/@storybook/addon-a11y": { - "version": "9.0.17", - "resolved": "https://registry.npmjs.org/@storybook/addon-a11y/-/addon-a11y-9.0.17.tgz", - "integrity": "sha512-9cXNK3q/atx3hwJAt9HkJbd9vUxCXfKKiNNuSACbf8h9/j6u3jktulKOf6Xjc3B8lwn6ZpdK/x1HHZN2kTqsvg==", + "version": "10.0.7", + "resolved": "https://registry.npmjs.org/@storybook/addon-a11y/-/addon-a11y-10.0.7.tgz", + "integrity": "sha512-JsYPpZ/n67/2bI1XJeyrAWHHQkHemPkPHjCA0tAUnMz1Shlo/LV2q1Ahgpxoihx4strbHwZz71bcS4MqkHBduA==", "dev": true, "license": "MIT", "dependencies": { @@ -1826,20 +1837,20 @@ "url": "https://opencollective.com/storybook" }, "peerDependencies": { - "storybook": "^9.0.17" + "storybook": "^10.0.7" } }, "node_modules/@storybook/addon-docs": { - "version": "9.0.17", - "resolved": "https://registry.npmjs.org/@storybook/addon-docs/-/addon-docs-9.0.17.tgz", - "integrity": "sha512-LOX/kKgQGnyulrqZHsvf77+ZoH/nSUaplGr5hvZglW/U6ak6fO9seJyXAzVKEnC6p+F8n02kFBZbi3s+znQhSg==", + "version": "10.0.7", + "resolved": "https://registry.npmjs.org/@storybook/addon-docs/-/addon-docs-10.0.7.tgz", + "integrity": "sha512-qQQMoeYZC4W+/8ubfOZiTrE8nYC/f4wWP1uq4peRyDy1N2nIN9SwhyxwMn0m3VpeGmRBga5dLvJY9ko6SnJekg==", "dev": true, "license": "MIT", "dependencies": { "@mdx-js/react": "^3.0.0", - "@storybook/csf-plugin": "9.0.17", - "@storybook/icons": "^1.2.12", - "@storybook/react-dom-shim": "9.0.17", + "@storybook/csf-plugin": "10.0.7", + "@storybook/icons": "^1.6.0", + "@storybook/react-dom-shim": "10.0.7", "react": "^16.8.0 || ^17.0.0 || ^18.0.0 || ^19.0.0", "react-dom": "^16.8.0 || ^17.0.0 || ^18.0.0 || ^19.0.0", "ts-dedent": "^2.0.0" @@ -1849,13 +1860,13 @@ "url": "https://opencollective.com/storybook" }, "peerDependencies": { - "storybook": "^9.0.17" + "storybook": "^10.0.7" } }, "node_modules/@storybook/addon-svelte-csf": { - "version": "5.0.7", - "resolved": "https://registry.npmjs.org/@storybook/addon-svelte-csf/-/addon-svelte-csf-5.0.7.tgz", - "integrity": "sha512-6Zmy5HjOlrrG6OoKRTGDr9LR6zRK4/Sa7raFzQRKHGASgMlfKsMdNTNO0sxnMUWCu2JMS6HsuoLtB3Ma8SlYtg==", + "version": "5.0.10", + "resolved": "https://registry.npmjs.org/@storybook/addon-svelte-csf/-/addon-svelte-csf-5.0.10.tgz", + "integrity": "sha512-poSvTS7VdaQ42ZoqW5e4+2Hv1iLO0mekH9fwn/QuBNse48R4WlTyR8XFbHRTfatl9gdc9ZYC4uWzazrmV6zGIA==", "dev": true, "license": "MIT", "dependencies": { @@ -1868,22 +1879,22 @@ "zimmerframe": "^1.1.2" }, "peerDependencies": { - "@storybook/svelte": "^0.0.0-0 || ^8.2.0 || ^9.0.0 || ^9.1.0-0", + "@storybook/svelte": "^0.0.0-0 || ^8.2.0 || ^9.0.0 || ^9.1.0-0 || ^10.0.0-0", "@sveltejs/vite-plugin-svelte": "^4.0.0 || ^5.0.0 || ^6.0.0", - "storybook": "^0.0.0-0 || ^8.2.0 || ^9.0.0 || ^9.1.0-0", + "storybook": "^0.0.0-0 || ^8.2.0 || ^9.0.0 || ^9.1.0-0 || ^10.0.0-0", "svelte": "^5.0.0", "vite": "^5.0.0 || ^6.0.0 || ^7.0.0" } }, "node_modules/@storybook/addon-vitest": { - "version": "9.0.17", - "resolved": "https://registry.npmjs.org/@storybook/addon-vitest/-/addon-vitest-9.0.17.tgz", - "integrity": "sha512-eogqcGbACR1sTedBSE2SP/4QV1ruicHYEhYjBtoPIjvYgymN1g5KSuQNysLx4f0SvAzczrcNjX2WVVLX2DVyzA==", + "version": "10.0.7", + "resolved": "https://registry.npmjs.org/@storybook/addon-vitest/-/addon-vitest-10.0.7.tgz", + "integrity": "sha512-i6v/mAl+elrUxb+1f4NdnM17t/fg+KGJWL1U9quflXTd3KiLY0xJB4LwNP6yYo7Imc5NIO2fRkJbGvNqLBRe2Q==", "dev": true, "license": "MIT", "dependencies": { "@storybook/global": "^5.0.0", - "@storybook/icons": "^1.4.0", + "@storybook/icons": "^1.6.0", "prompts": "^2.4.0", "ts-dedent": "^2.2.0" }, @@ -1892,15 +1903,19 @@ "url": "https://opencollective.com/storybook" }, "peerDependencies": { - "@vitest/browser": "^3.0.0", - "@vitest/runner": "^3.0.0", - "storybook": "^9.0.17", - "vitest": "^3.0.0" + "@vitest/browser": "^3.0.0 || ^4.0.0", + "@vitest/browser-playwright": "^4.0.0", + "@vitest/runner": "^3.0.0 || ^4.0.0", + "storybook": "^10.0.7", + "vitest": "^3.0.0 || ^4.0.0" }, "peerDependenciesMeta": { "@vitest/browser": { "optional": true }, + "@vitest/browser-playwright": { + "optional": true + }, "@vitest/runner": { "optional": true }, @@ -1910,13 +1925,13 @@ } }, "node_modules/@storybook/builder-vite": { - "version": "9.0.17", - "resolved": "https://registry.npmjs.org/@storybook/builder-vite/-/builder-vite-9.0.17.tgz", - "integrity": "sha512-lyuvgGhb0NaVk1tdB4xwzky6+YXQfxlxfNQqENYZ9uYQZdPfErMa4ZTXVQTV+CQHAa2NL+p/dG2JPAeu39e9UA==", + "version": "10.0.7", + "resolved": "https://registry.npmjs.org/@storybook/builder-vite/-/builder-vite-10.0.7.tgz", + "integrity": "sha512-wk2TAoUY5+9t78GWVBndu9rEo9lo6Ec3SRrLT4VpIlcS2GPK+5f26UC2uvIBwOF/N7JrUUKq/zWDZ3m+do9QDg==", "dev": true, "license": "MIT", "dependencies": { - "@storybook/csf-plugin": "9.0.17", + "@storybook/csf-plugin": "10.0.7", "ts-dedent": "^2.0.0" }, "funding": { @@ -1924,7 +1939,7 @@ "url": "https://opencollective.com/storybook" }, "peerDependencies": { - "storybook": "^9.0.17", + "storybook": "^10.0.7", "vite": "^5.0.0 || ^6.0.0 || ^7.0.0" } }, @@ -1939,20 +1954,38 @@ } }, "node_modules/@storybook/csf-plugin": { - "version": "9.0.17", - "resolved": "https://registry.npmjs.org/@storybook/csf-plugin/-/csf-plugin-9.0.17.tgz", - "integrity": "sha512-6Q4eo1ObrLlsnB6bIt6T8+45XAb4to2pQGNrI7QPkLQRLrZinrJcNbLY7AGkyIoCOEsEbq08n09/nClQUbu8HA==", + "version": "10.0.7", + "resolved": "https://registry.npmjs.org/@storybook/csf-plugin/-/csf-plugin-10.0.7.tgz", + "integrity": "sha512-YaYYlCyJBwxaMk7yREOdz+9MDSgxIYGdeJ9EIq/bUndmkoj9SRo1P9/0lC5dseWQoiGy4T3PbZiWruD8uM5m3g==", "dev": true, "license": "MIT", "dependencies": { - "unplugin": "^1.3.1" + "unplugin": "^2.3.5" }, "funding": { "type": "opencollective", "url": "https://opencollective.com/storybook" }, "peerDependencies": { - "storybook": "^9.0.17" + "esbuild": "*", + "rollup": "*", + "storybook": "^10.0.7", + "vite": "*", + "webpack": "*" + }, + "peerDependenciesMeta": { + "esbuild": { + "optional": true + }, + "rollup": { + "optional": true + }, + "vite": { + "optional": true + }, + "webpack": { + "optional": true + } } }, "node_modules/@storybook/global": { @@ -1963,9 +1996,9 @@ "license": "MIT" }, "node_modules/@storybook/icons": { - "version": "1.4.0", - "resolved": "https://registry.npmjs.org/@storybook/icons/-/icons-1.4.0.tgz", - "integrity": "sha512-Td73IeJxOyalzvjQL+JXx72jlIYHgs+REaHiREOqfpo3A2AYYG71AUbcv+lg7mEDIweKVCxsMQ0UKo634c8XeA==", + "version": "1.6.0", + "resolved": "https://registry.npmjs.org/@storybook/icons/-/icons-1.6.0.tgz", + "integrity": "sha512-hcFZIjW8yQz8O8//2WTIXylm5Xsgc+lW9ISLgUk1xGmptIJQRdlhVIXCpSyLrQaaRiyhQRaVg7l3BD9S216BHw==", "dev": true, "license": "MIT", "engines": { @@ -1977,9 +2010,9 @@ } }, "node_modules/@storybook/react-dom-shim": { - "version": "9.0.17", - "resolved": "https://registry.npmjs.org/@storybook/react-dom-shim/-/react-dom-shim-9.0.17.tgz", - "integrity": "sha512-ak/x/m6MDDxdE6rCDymTltaiQF3oiKrPHSwfM+YPgQR6MVmzTTs4+qaPfeev7FZEHq23IkfDMTmSTTJtX7Vs9A==", + "version": "10.0.7", + "resolved": "https://registry.npmjs.org/@storybook/react-dom-shim/-/react-dom-shim-10.0.7.tgz", + "integrity": "sha512-bp4OnMtZGwPJQDqNRi4K5iibLbZ2TZZMkWW7oSw5jjPFpGSreSjCe8LH9yj/lDnK8Ox9bGMCBFE5RV5XuML29w==", "dev": true, "license": "MIT", "funding": { @@ -1987,126 +2020,75 @@ "url": "https://opencollective.com/storybook" }, "peerDependencies": { - "react": "^16.8.0 || ^17.0.0 || ^18.0.0 || ^19.0.0-beta", - "react-dom": "^16.8.0 || ^17.0.0 || ^18.0.0 || ^19.0.0-beta", - "storybook": "^9.0.17" + "react": "^16.8.0 || ^17.0.0 || ^18.0.0 || ^19.0.0", + "react-dom": "^16.8.0 || ^17.0.0 || ^18.0.0 || ^19.0.0", + "storybook": "^10.0.7" } }, "node_modules/@storybook/svelte": { - "version": "9.0.17", - "resolved": "https://registry.npmjs.org/@storybook/svelte/-/svelte-9.0.17.tgz", - "integrity": "sha512-RwOswdq7S3+ZOuoM/oRrcmlsKdjcd/3wMHbuirzYoAhdwsjubSuRepMV64O9RnlXd3x7rZw4fXpq1M/SVo5XiQ==", + "version": "10.0.7", + "resolved": "https://registry.npmjs.org/@storybook/svelte/-/svelte-10.0.7.tgz", + "integrity": "sha512-rO+YQhHucy47Vh67z318pALmd6x+K1Kj30Fb4a6oOEw4xn4zCo9KTmkMWs24c4oduEXD/eJu3badlRmsVXzyfA==", "dev": true, "license": "MIT", "dependencies": { "ts-dedent": "^2.0.0", "type-fest": "~2.19" }, - "engines": { - "node": ">=20.0.0" - }, "funding": { "type": "opencollective", "url": "https://opencollective.com/storybook" }, "peerDependencies": { - "storybook": "^9.0.17", + "storybook": "^10.0.7", "svelte": "^5.0.0" } }, - "node_modules/@storybook/sveltekit": { - "version": "9.0.17", - "resolved": "https://registry.npmjs.org/@storybook/sveltekit/-/sveltekit-9.0.17.tgz", - "integrity": "sha512-CUOATuW5Qk3SjNvmjH+wyx2GCsMF1cvw3gwkujV9kehPebzV20NhgHpbzSoepvwF7+Bj6jl8V6UxiMWk0jJFmA==", + "node_modules/@storybook/svelte-vite": { + "version": "10.0.7", + "resolved": "https://registry.npmjs.org/@storybook/svelte-vite/-/svelte-vite-10.0.7.tgz", + "integrity": "sha512-q9/RtrhX1CnznO6AO9MDEy1bsccbGeRxW28FLpgUrztV4IGZ/dFUrFIFurKRyuA3/nFsbtzp1F5jFt3RExmmTw==", "dev": true, "license": "MIT", "dependencies": { - "@storybook/builder-vite": "9.0.17", - "@storybook/svelte": "9.0.17", - "@storybook/svelte-vite": "9.0.17" - }, - "engines": { - "node": ">=20.0.0" - }, - "funding": { - "type": "opencollective", - "url": "https://opencollective.com/storybook" - }, - "peerDependencies": { - "storybook": "^9.0.17", - "svelte": "^5.0.0", - "vite": "^5.0.0 || ^6.0.0 || ^7.0.0" - } - }, - "node_modules/@storybook/sveltekit/node_modules/@storybook/svelte-vite": { - "version": "9.0.17", - "resolved": "https://registry.npmjs.org/@storybook/svelte-vite/-/svelte-vite-9.0.17.tgz", - "integrity": "sha512-fRIxOZy9IRI6BfL1LgFn+B+IckGOlT1SstD01y9ddO4pVKWih/l+vb44bnZs+Z0faJZbrG/LgfnXTOPj052Z8g==", - "dev": true, - "license": "MIT", - "dependencies": { - "@storybook/builder-vite": "9.0.17", - "@storybook/svelte": "9.0.17", + "@storybook/builder-vite": "10.0.7", + "@storybook/svelte": "10.0.7", "magic-string": "^0.30.0", - "svelte2tsx": "^0.7.35", + "svelte2tsx": "^0.7.44", "typescript": "^4.9.4 || ^5.0.0" }, - "engines": { - "node": ">=20.0.0" + "funding": { + "type": "opencollective", + "url": "https://opencollective.com/storybook" + }, + "peerDependencies": { + "@sveltejs/vite-plugin-svelte": "^2.0.0 || ^3.0.0 || ^4.0.0 || ^5.0.0 || ^6.0.0", + "storybook": "^10.0.7", + "svelte": "^5.0.0", + "vite": "^5.0.0 || ^6.0.0 || ^7.0.0" + } + }, + "node_modules/@storybook/sveltekit": { + "version": "10.0.7", + "resolved": "https://registry.npmjs.org/@storybook/sveltekit/-/sveltekit-10.0.7.tgz", + "integrity": "sha512-ujTW7PfWvgBrzd7jzaZe9JgjUeM5YvBKm+xru6t7Dr4bdfmkKqlZHPRdXn/sy+fQNyfg6JL2WKy2KIIeA+RvSg==", + "dev": true, + "license": "MIT", + "dependencies": { + "@storybook/builder-vite": "10.0.7", + "@storybook/svelte": "10.0.7", + "@storybook/svelte-vite": "10.0.7" }, "funding": { "type": "opencollective", "url": "https://opencollective.com/storybook" }, "peerDependencies": { - "@sveltejs/vite-plugin-svelte": "^2.0.0 || ^3.0.0 || ^4.0.0 || ^5.0.0", - "storybook": "^9.0.17", + "storybook": "^10.0.7", "svelte": "^5.0.0", "vite": "^5.0.0 || ^6.0.0 || ^7.0.0" } }, - "node_modules/@storybook/sveltekit/node_modules/@sveltejs/vite-plugin-svelte": { - "version": "5.1.1", - "resolved": "https://registry.npmjs.org/@sveltejs/vite-plugin-svelte/-/vite-plugin-svelte-5.1.1.tgz", - "integrity": "sha512-Y1Cs7hhTc+a5E9Va/xwKlAJoariQyHY+5zBgCZg4PFWNYQ1nMN9sjK1zhw1gK69DuqVP++sht/1GZg1aRwmAXQ==", - "dev": true, - "license": "MIT", - "peer": true, - "dependencies": { - "@sveltejs/vite-plugin-svelte-inspector": "^4.0.1", - "debug": "^4.4.1", - "deepmerge": "^4.3.1", - "kleur": "^4.1.5", - "magic-string": "^0.30.17", - "vitefu": "^1.0.6" - }, - "engines": { - "node": "^18.0.0 || ^20.0.0 || >=22" - }, - "peerDependencies": { - "svelte": "^5.0.0", - "vite": "^6.0.0" - } - }, - "node_modules/@storybook/sveltekit/node_modules/@sveltejs/vite-plugin-svelte/node_modules/@sveltejs/vite-plugin-svelte-inspector": { - "version": "4.0.1", - "resolved": "https://registry.npmjs.org/@sveltejs/vite-plugin-svelte-inspector/-/vite-plugin-svelte-inspector-4.0.1.tgz", - "integrity": "sha512-J/Nmb2Q2y7mck2hyCX4ckVHcR5tu2J+MtBEQqpDrrgELZ2uvraQcK/ioCV61AqkdXFgriksOKIceDcQmqnGhVw==", - "dev": true, - "license": "MIT", - "peer": true, - "dependencies": { - "debug": "^4.3.7" - }, - "engines": { - "node": "^18.0.0 || ^20.0.0 || >=22" - }, - "peerDependencies": { - "@sveltejs/vite-plugin-svelte": "^5.0.0", - "svelte": "^5.0.0", - "vite": "^6.0.0" - } - }, "node_modules/@sveltejs/acorn-typescript": { "version": "1.0.5", "resolved": "https://registry.npmjs.org/@sveltejs/acorn-typescript/-/acorn-typescript-1.0.5.tgz", @@ -2117,9 +2099,9 @@ } }, "node_modules/@sveltejs/adapter-static": { - "version": "3.0.9", - "resolved": "https://registry.npmjs.org/@sveltejs/adapter-static/-/adapter-static-3.0.9.tgz", - "integrity": "sha512-aytHXcMi7lb9ljsWUzXYQ0p5X1z9oWud2olu/EpmH7aCu4m84h7QLvb5Wp+CFirKcwoNnYvYWhyP/L8Vh1ztdw==", + "version": "3.0.10", + "resolved": "https://registry.npmjs.org/@sveltejs/adapter-static/-/adapter-static-3.0.10.tgz", + "integrity": "sha512-7D9lYFWJmB7zxZyTE/qxjksvMqzMuYrrsyh1f4AlZqeZeACPRySjbC3aFiY55wb1tWUaKOQG9PVbm74JcN2Iew==", "dev": true, "license": "MIT", "peerDependencies": { @@ -2127,9 +2109,9 @@ } }, "node_modules/@sveltejs/kit": { - "version": "2.37.0", - "resolved": "https://registry.npmjs.org/@sveltejs/kit/-/kit-2.37.0.tgz", - "integrity": "sha512-xgKtpjQ6Ry4mdShd01ht5AODUsW7+K1iValPDq7QX8zI1hWOKREH9GjG8SRCN5tC4K7UXmMhuQam7gbLByVcnw==", + "version": "2.48.4", + "resolved": "https://registry.npmjs.org/@sveltejs/kit/-/kit-2.48.4.tgz", + "integrity": "sha512-TGFX1pZUt9qqY20Cv5NyYvy0iLWHf2jXi8s+eCGsig7jQMdwZWKUFMR6TbvFNhfDSUpc1sH/Y5EHv20g3HHA3g==", "dev": true, "license": "MIT", "dependencies": { @@ -2166,16 +2148,15 @@ } }, "node_modules/@sveltejs/vite-plugin-svelte": { - "version": "6.1.0", - "resolved": "https://registry.npmjs.org/@sveltejs/vite-plugin-svelte/-/vite-plugin-svelte-6.1.0.tgz", - "integrity": "sha512-+U6lz1wvGEG/BvQyL4z/flyNdQ9xDNv5vrh+vWBWTHaebqT0c9RNggpZTo/XSPoHsSCWBlYaTlRX8pZ9GATXCw==", + "version": "6.2.1", + "resolved": "https://registry.npmjs.org/@sveltejs/vite-plugin-svelte/-/vite-plugin-svelte-6.2.1.tgz", + "integrity": "sha512-YZs/OSKOQAQCnJvM/P+F1URotNnYNeU3P2s4oIpzm1uFaqUEqRxUB0g5ejMjEb5Gjb9/PiBI5Ktrq4rUUF8UVQ==", "dev": true, "license": "MIT", "dependencies": { - "@sveltejs/vite-plugin-svelte-inspector": "^5.0.0-next.1", + "@sveltejs/vite-plugin-svelte-inspector": "^5.0.0", "debug": "^4.4.1", "deepmerge": "^4.3.1", - "kleur": "^4.1.5", "magic-string": "^0.30.17", "vitefu": "^1.1.1" }, @@ -3361,19 +3342,6 @@ "node": ">= 0.8" } }, - "node_modules/better-opn": { - "version": "3.0.2", - "resolved": "https://registry.npmjs.org/better-opn/-/better-opn-3.0.2.tgz", - "integrity": "sha512-aVNobHnJqLiUelTaHat9DZ1qM2w0C0Eym4LPI/3JxOnSokGVdsl1T1kN7TFvsEAD8G47A6VKQ0TVHqbBnYMJlQ==", - "dev": true, - "license": "MIT", - "dependencies": { - "open": "^8.0.4" - }, - "engines": { - "node": ">=12.0.0" - } - }, "node_modules/bits-ui": { "version": "2.8.11", "resolved": "https://registry.npmjs.org/bits-ui/-/bits-ui-2.8.11.tgz", @@ -3844,16 +3812,6 @@ "node": ">=0.10.0" } }, - "node_modules/define-lazy-prop": { - "version": "2.0.0", - "resolved": "https://registry.npmjs.org/define-lazy-prop/-/define-lazy-prop-2.0.0.tgz", - "integrity": "sha512-Ds09qNh8yw3khSjiJjiUInaGX9xlqZDY7JVryGxdxV7NPeuqQfplOpQ66yJFZut3jLa5zOwkXw1g9EI2uKh4Og==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=8" - } - }, "node_modules/dequal": { "version": "2.0.3", "resolved": "https://registry.npmjs.org/dequal/-/dequal-2.0.3.tgz", @@ -4042,19 +4000,6 @@ "@esbuild/win32-x64": "0.25.8" } }, - "node_modules/esbuild-register": { - "version": "3.6.0", - "resolved": "https://registry.npmjs.org/esbuild-register/-/esbuild-register-3.6.0.tgz", - "integrity": "sha512-H2/S7Pm8a9CL1uhp9OvjwrBh5Pvx0H8qVOxNu8Wed9Y7qv56MPtq+GGM8RJpq6glYJn9Wspr8uw7l55uyinNeg==", - "dev": true, - "license": "MIT", - "dependencies": { - "debug": "^4.3.4" - }, - "peerDependencies": { - "esbuild": ">=0.12 <1" - } - }, "node_modules/escape-string-regexp": { "version": "4.0.0", "resolved": "https://registry.npmjs.org/escape-string-regexp/-/escape-string-regexp-4.0.0.tgz", @@ -4146,20 +4091,17 @@ } }, "node_modules/eslint-plugin-storybook": { - "version": "9.0.17", - "resolved": "https://registry.npmjs.org/eslint-plugin-storybook/-/eslint-plugin-storybook-9.0.17.tgz", - "integrity": "sha512-IuTdlwCEwoDNobdygRCxNhlKXHmsDfPtPvHGcsY35x2Bx8KItrjfekO19gJrjc1VT2CMfcZMYF8OBKaxHELupw==", + "version": "10.0.7", + "resolved": "https://registry.npmjs.org/eslint-plugin-storybook/-/eslint-plugin-storybook-10.0.7.tgz", + "integrity": "sha512-qOQq9KdT1jsBgT3qsxUH2n67aj1WR8D1XCoER8Q6yuVlS5TimNwk1mZeWkXVf/o4RQQT6flT2y5cG2gPLZPvJA==", "dev": true, "license": "MIT", "dependencies": { "@typescript-eslint/utils": "^8.8.1" }, - "engines": { - "node": ">=20.0.0" - }, "peerDependencies": { "eslint": ">=8", - "storybook": "^9.0.17" + "storybook": "^10.0.7" } }, "node_modules/eslint-plugin-svelte": { @@ -4405,11 +4347,14 @@ } }, "node_modules/fdir": { - "version": "6.4.6", - "resolved": "https://registry.npmjs.org/fdir/-/fdir-6.4.6.tgz", - "integrity": "sha512-hiFoqpyZcfNm1yc4u8oWCf9A2c4D3QjCrks3zmoVKVxpQRzmPNar1hUJcBG2RQHvEVGDN+Jm81ZheVLAQMK6+w==", + "version": "6.5.0", + "resolved": "https://registry.npmjs.org/fdir/-/fdir-6.5.0.tgz", + "integrity": "sha512-tIbYtZbucOs0BRGqPJkshJUYdL+SDH7dVM8gjy+ERp3WAUjLEFJE+02kanyHtwjWOnwrKYBiwAmM0p4kLJAnXg==", "dev": true, "license": "MIT", + "engines": { + "node": ">=12.0.0" + }, "peerDependencies": { "picomatch": "^3 || ^4" }, @@ -5072,22 +5017,6 @@ "integrity": "sha512-0aO8FkhNZlj/ZIbNi7Lxxr12obT7cL1moPfE4tg1LkX7LlLfC6DeX4l2ZEud1ukP9jNQyNnfzQVqwbwmAATY4Q==", "license": "MIT" }, - "node_modules/is-docker": { - "version": "2.2.1", - "resolved": "https://registry.npmjs.org/is-docker/-/is-docker-2.2.1.tgz", - "integrity": "sha512-F+i2BKsFrH66iaUFc0woD8sLy8getkwTwtOBjvs56Cx4CgJDeKQeqfz8wAYiSb8JOprWhHH5p77PbmYCvvUuXQ==", - "dev": true, - "license": "MIT", - "bin": { - "is-docker": "cli.js" - }, - "engines": { - "node": ">=8" - }, - "funding": { - "url": "https://github.com/sponsors/sindresorhus" - } - }, "node_modules/is-extglob": { "version": "2.1.1", "resolved": "https://registry.npmjs.org/is-extglob/-/is-extglob-2.1.1.tgz", @@ -5133,19 +5062,6 @@ "url": "https://github.com/sponsors/sindresorhus" } }, - "node_modules/is-wsl": { - "version": "2.2.0", - "resolved": "https://registry.npmjs.org/is-wsl/-/is-wsl-2.2.0.tgz", - "integrity": "sha512-fKzAra0rGJUUBwGBgNkHZuToZcn+TtXHpeCgmkMJMMYx1sQDYaCSyjJBSCa2nH1DGm7s3n1oBnohoVTBaN7Lww==", - "dev": true, - "license": "MIT", - "dependencies": { - "is-docker": "^2.0.0" - }, - "engines": { - "node": ">=8" - } - }, "node_modules/isexe": { "version": "2.0.0", "resolved": "https://registry.npmjs.org/isexe/-/isexe-2.0.0.tgz", @@ -5591,16 +5507,6 @@ "dev": true, "license": "MIT" }, - "node_modules/lower-case": { - "version": "2.0.2", - "resolved": "https://registry.npmjs.org/lower-case/-/lower-case-2.0.2.tgz", - "integrity": "sha512-7fm3l3NAF9WfN6W3JOmf5drwpVqX78JtoGJ3A6W0a6ZnldM41w2fV5D490psKFTpMds8TJse/eHLFFsNHHjHgg==", - "dev": true, - "license": "MIT", - "dependencies": { - "tslib": "^2.0.3" - } - }, "node_modules/lowlight": { "version": "3.3.0", "resolved": "https://registry.npmjs.org/lowlight/-/lowlight-3.3.0.tgz", @@ -6783,17 +6689,6 @@ "dev": true, "license": "MIT" }, - "node_modules/no-case": { - "version": "3.0.4", - "resolved": "https://registry.npmjs.org/no-case/-/no-case-3.0.4.tgz", - "integrity": "sha512-fgAN3jGAh+RoxUGZHTSOLJIqUc2wmoBwGR4tbpNAKmmovFoWq0OdRkb0VkldReO2a2iBT/OEulG9XSUc10r3zg==", - "dev": true, - "license": "MIT", - "dependencies": { - "lower-case": "^2.0.2", - "tslib": "^2.0.3" - } - }, "node_modules/node-addon-api": { "version": "7.1.1", "resolved": "https://registry.npmjs.org/node-addon-api/-/node-addon-api-7.1.1.tgz", @@ -6815,24 +6710,6 @@ "url": "https://github.com/sponsors/ljharb" } }, - "node_modules/open": { - "version": "8.4.2", - "resolved": "https://registry.npmjs.org/open/-/open-8.4.2.tgz", - "integrity": "sha512-7x81NCL719oNbsq/3mh+hVrAWmFuEYUqrq/Iw3kUzH8ReypT9QQ0BLoJS7/G9k6N81XjW4qHWtjWwe/9eLy1EQ==", - "dev": true, - "license": "MIT", - "dependencies": { - "define-lazy-prop": "^2.0.0", - "is-docker": "^2.1.1", - "is-wsl": "^2.2.0" - }, - "engines": { - "node": ">=12" - }, - "funding": { - "url": "https://github.com/sponsors/sindresorhus" - } - }, "node_modules/opener": { "version": "1.5.2", "resolved": "https://registry.npmjs.org/opener/-/opener-1.5.2.tgz", @@ -6919,17 +6796,6 @@ "url": "https://github.com/inikulin/parse5?sponsor=1" } }, - "node_modules/pascal-case": { - "version": "3.1.2", - "resolved": "https://registry.npmjs.org/pascal-case/-/pascal-case-3.1.2.tgz", - "integrity": "sha512-uWlGT3YSnK9x3BQJaOdcZwrnV6hPpd8jFH1/ucpiLRPh/2zCVJKS19E4GvYHvaCcACn3foXZ0cLB9Wrx1KGe5g==", - "dev": true, - "license": "MIT", - "dependencies": { - "no-case": "^3.0.4", - "tslib": "^2.0.3" - } - }, "node_modules/path-exists": { "version": "4.0.0", "resolved": "https://registry.npmjs.org/path-exists/-/path-exists-4.0.0.tgz", @@ -7000,13 +6866,13 @@ } }, "node_modules/playwright": { - "version": "1.54.1", - "resolved": "https://registry.npmjs.org/playwright/-/playwright-1.54.1.tgz", - "integrity": "sha512-peWpSwIBmSLi6aW2auvrUtf2DqY16YYcCMO8rTVx486jKmDTJg7UAhyrraP98GB8BoPURZP8+nxO7TSd4cPr5g==", + "version": "1.56.1", + "resolved": "https://registry.npmjs.org/playwright/-/playwright-1.56.1.tgz", + "integrity": "sha512-aFi5B0WovBHTEvpM3DzXTUaeN6eN0qWnTkKx4NQaH4Wvcmc153PdaY2UBdSYKaGYw+UyWXSVyxDUg5DoPEttjw==", "dev": true, "license": "Apache-2.0", "dependencies": { - "playwright-core": "1.54.1" + "playwright-core": "1.56.1" }, "bin": { "playwright": "cli.js" @@ -7019,9 +6885,9 @@ } }, "node_modules/playwright-core": { - "version": "1.54.1", - "resolved": "https://registry.npmjs.org/playwright-core/-/playwright-core-1.54.1.tgz", - "integrity": "sha512-Nbjs2zjj0htNhzgiy5wu+3w09YetDx5pkrpI/kZotDlDUaYk0HVA5xrBVPdow4SAUIlhgKcJeJg4GRKW6xHusA==", + "version": "1.56.1", + "resolved": "https://registry.npmjs.org/playwright-core/-/playwright-core-1.56.1.tgz", + "integrity": "sha512-hutraynyn31F+Bifme+Ps9Vq59hKuUCz7H1kDOcBs+2oGguKkWTU50bBWrtz34OUWmIwpBTWDxaRPXrIXkgvmQ==", "dev": true, "license": "Apache-2.0", "bin": { @@ -7852,6 +7718,13 @@ "dev": true, "license": "MIT" }, + "node_modules/scule": { + "version": "1.3.0", + "resolved": "https://registry.npmjs.org/scule/-/scule-1.3.0.tgz", + "integrity": "sha512-6FtHJEvt+pVMIB9IBY+IcCJ6Z5f1iQnytgyfKMhDKgmzYG+TeH/wx1y3l27rshSbLiSanrR9ffZDrEsmjlQF2g==", + "dev": true, + "license": "MIT" + }, "node_modules/secure-compare": { "version": "3.0.1", "resolved": "https://registry.npmjs.org/secure-compare/-/secure-compare-3.0.1.tgz", @@ -8052,26 +7925,26 @@ "license": "MIT" }, "node_modules/storybook": { - "version": "9.0.17", - "resolved": "https://registry.npmjs.org/storybook/-/storybook-9.0.17.tgz", - "integrity": "sha512-O+9jgJ+Trlq9VGD1uY4OBLKQWHHDKM/A/pA8vMW6PVehhGHNvpzcIC1bngr6mL5gGHZP2nBv+9XG8pTMcggMmg==", + "version": "10.0.7", + "resolved": "https://registry.npmjs.org/storybook/-/storybook-10.0.7.tgz", + "integrity": "sha512-7smAu0o+kdm378Q2uIddk32pn0UdIbrtTVU+rXRVtTVTCrK/P2cCui2y4JH+Bl3NgEq1bbBQpCAF/HKrDjk2Qw==", "dev": true, "license": "MIT", "dependencies": { "@storybook/global": "^5.0.0", + "@storybook/icons": "^1.6.0", "@testing-library/jest-dom": "^6.6.3", "@testing-library/user-event": "^14.6.1", "@vitest/expect": "3.2.4", + "@vitest/mocker": "3.2.4", "@vitest/spy": "3.2.4", - "better-opn": "^3.0.2", "esbuild": "^0.18.0 || ^0.19.0 || ^0.20.0 || ^0.21.0 || ^0.22.0 || ^0.23.0 || ^0.24.0 || ^0.25.0", - "esbuild-register": "^3.5.0", "recast": "^0.23.5", "semver": "^7.6.2", "ws": "^8.18.0" }, "bin": { - "storybook": "bin/index.cjs" + "storybook": "dist/bin/dispatcher.js" }, "funding": { "type": "opencollective", @@ -8418,14 +8291,14 @@ } }, "node_modules/svelte2tsx": { - "version": "0.7.41", - "resolved": "https://registry.npmjs.org/svelte2tsx/-/svelte2tsx-0.7.41.tgz", - "integrity": "sha512-/TUwpyn/Qc1wcGuayf2GSwvZ7htdAOzpo0JFFm96srKnRXoTD0gy4n06g+XgH8w016S3lPtyFVtFAm+0yJ0BZw==", + "version": "0.7.45", + "resolved": "https://registry.npmjs.org/svelte2tsx/-/svelte2tsx-0.7.45.tgz", + "integrity": "sha512-cSci+mYGygYBHIZLHlm/jYlEc1acjAHqaQaDFHdEBpUueM9kSTnPpvPtSl5VkJOU1qSJ7h1K+6F/LIUYiqC8VA==", "dev": true, "license": "MIT", "dependencies": { "dedent-js": "^1.0.1", - "pascal-case": "^3.1.1" + "scule": "^1.3.0" }, "peerDependencies": { "svelte": "^3.55 || ^4.0.0-next.0 || ^4.0 || ^5.0.0-next.0", @@ -8535,14 +8408,14 @@ "license": "MIT" }, "node_modules/tinyglobby": { - "version": "0.2.14", - "resolved": "https://registry.npmjs.org/tinyglobby/-/tinyglobby-0.2.14.tgz", - "integrity": "sha512-tX5e7OM1HnYr2+a2C/4V0htOcSQcoSTH9KgJnVvNm5zm/cyEWKJ7j7YutsH9CxMdtOkkLFy2AHrMci9IM8IPZQ==", + "version": "0.2.15", + "resolved": "https://registry.npmjs.org/tinyglobby/-/tinyglobby-0.2.15.tgz", + "integrity": "sha512-j2Zq4NyQYG5XMST4cbs02Ak8iJUdxRM0XI5QyxXuZOzKOINmWurp3smXu3y5wDcJrptwpSjgXHzIQxR0omXljQ==", "dev": true, "license": "MIT", "dependencies": { - "fdir": "^6.4.4", - "picomatch": "^4.0.2" + "fdir": "^6.5.0", + "picomatch": "^4.0.3" }, "engines": { "node": ">=12.0.0" @@ -8918,17 +8791,19 @@ } }, "node_modules/unplugin": { - "version": "1.16.1", - "resolved": "https://registry.npmjs.org/unplugin/-/unplugin-1.16.1.tgz", - "integrity": "sha512-4/u/j4FrCKdi17jaxuJA0jClGxB1AvU2hw/IuayPc4ay1XGaJs/rbb4v5WKwAjNifjmXK9PIFyuPiaK8azyR9w==", + "version": "2.3.10", + "resolved": "https://registry.npmjs.org/unplugin/-/unplugin-2.3.10.tgz", + "integrity": "sha512-6NCPkv1ClwH+/BGE9QeoTIl09nuiAt0gS28nn1PvYXsGKRwM2TCbFA2QiilmehPDTXIe684k4rZI1yl3A1PCUw==", "dev": true, "license": "MIT", "dependencies": { - "acorn": "^8.14.0", + "@jridgewell/remapping": "^2.3.5", + "acorn": "^8.15.0", + "picomatch": "^4.0.3", "webpack-virtual-modules": "^0.6.2" }, "engines": { - "node": ">=14.0.0" + "node": ">=18.12.0" } }, "node_modules/uri-js": { @@ -9054,18 +8929,18 @@ } }, "node_modules/vite": { - "version": "7.0.5", - "resolved": "https://registry.npmjs.org/vite/-/vite-7.0.5.tgz", - "integrity": "sha512-1mncVwJxy2C9ThLwz0+2GKZyEXuC3MyWtAAlNftlZZXZDP3AJt5FmwcMit/IGGaNZ8ZOB2BNO/HFUB+CpN0NQw==", + "version": "7.2.2", + "resolved": "https://registry.npmjs.org/vite/-/vite-7.2.2.tgz", + "integrity": "sha512-BxAKBWmIbrDgrokdGZH1IgkIk/5mMHDreLDmCJ0qpyJaAteP8NvMhkwr/ZCQNqNH97bw/dANTE9PDzqwJghfMQ==", "dev": true, "license": "MIT", "dependencies": { "esbuild": "^0.25.0", - "fdir": "^6.4.6", - "picomatch": "^4.0.2", + "fdir": "^6.5.0", + "picomatch": "^4.0.3", "postcss": "^8.5.6", - "rollup": "^4.40.0", - "tinyglobby": "^0.2.14" + "rollup": "^4.43.0", + "tinyglobby": "^0.2.15" }, "bin": { "vite": "bin/vite.js" diff --git a/tools/server/webui/package.json b/tools/server/webui/package.json index 92c7457bd..8b88f691a 100644 --- a/tools/server/webui/package.json +++ b/tools/server/webui/package.json @@ -24,20 +24,20 @@ "cleanup": "rm -rf .svelte-kit build node_modules test-results" }, "devDependencies": { - "@chromatic-com/storybook": "^4.0.1", + "@chromatic-com/storybook": "^4.1.2", "@eslint/compat": "^1.2.5", "@eslint/js": "^9.18.0", "@internationalized/date": "^3.8.2", "@lucide/svelte": "^0.515.0", "@playwright/test": "^1.49.1", - "@storybook/addon-a11y": "^9.0.17", - "@storybook/addon-docs": "^9.0.17", - "@storybook/addon-svelte-csf": "^5.0.7", - "@storybook/addon-vitest": "^9.0.17", - "@storybook/sveltekit": "^9.0.17", - "@sveltejs/adapter-static": "^3.0.8", - "@sveltejs/kit": "^2.22.0", - "@sveltejs/vite-plugin-svelte": "^6.0.0", + "@storybook/addon-a11y": "^10.0.7", + "@storybook/addon-docs": "^10.0.7", + "@storybook/addon-svelte-csf": "^5.0.10", + "@storybook/addon-vitest": "^10.0.7", + "@storybook/sveltekit": "^10.0.7", + "@sveltejs/adapter-static": "^3.0.10", + "@sveltejs/kit": "^2.48.4", + "@sveltejs/vite-plugin-svelte": "^6.2.1", "@tailwindcss/forms": "^0.5.9", "@tailwindcss/typography": "^0.5.15", "@tailwindcss/vite": "^4.0.0", @@ -48,21 +48,21 @@ "dexie": "^4.0.11", "eslint": "^9.18.0", "eslint-config-prettier": "^10.0.1", - "eslint-plugin-storybook": "^9.0.17", + "eslint-plugin-storybook": "^10.0.7", "eslint-plugin-svelte": "^3.0.0", "fflate": "^0.8.2", "globals": "^16.0.0", "http-server": "^14.1.1", "mdast": "^3.0.0", "mdsvex": "^0.12.3", - "playwright": "^1.53.0", + "playwright": "^1.56.1", "prettier": "^3.4.2", "prettier-plugin-svelte": "^3.3.3", "prettier-plugin-tailwindcss": "^0.6.11", "rehype-katex": "^7.0.1", "remark-math": "^6.0.0", "sass": "^1.93.3", - "storybook": "^9.0.17", + "storybook": "^10.0.7", "svelte": "^5.0.0", "svelte-check": "^4.0.0", "tailwind-merge": "^3.3.1", @@ -73,7 +73,7 @@ "typescript-eslint": "^8.20.0", "unified": "^11.0.5", "uuid": "^13.0.0", - "vite": "^7.0.4", + "vite": "^7.2.2", "vite-plugin-devtools-json": "^0.2.0", "vitest": "^3.2.3", "vitest-browser-svelte": "^0.1.0" diff --git a/tools/server/webui/src/stories/ChatForm.stories.svelte b/tools/server/webui/src/stories/ChatForm.stories.svelte index 6a0fc0870..82848e4fb 100644 --- a/tools/server/webui/src/stories/ChatForm.stories.svelte +++ b/tools/server/webui/src/stories/ChatForm.stories.svelte @@ -1,7 +1,7 @@