From 0e1d9185c5fe82e905d1f5ae6b2e5dcd607a8dfd Mon Sep 17 00:00:00 2001 From: Eve <139727413+netrunnereve@users.noreply.github.com> Date: Thu, 20 Aug 2026 20:01:32 +0000 Subject: [PATCH 01/38] ci: use shell script to check cmake pkg (#27414) * use regular script to build cmake pkg * use old grep without perl --- .github/workflows/build-cmake-pkg.yml | 36 ++++++++++++--------------- 1 file changed, 16 insertions(+), 20 deletions(-) diff --git a/.github/workflows/build-cmake-pkg.yml b/.github/workflows/build-cmake-pkg.yml index 0e4069ce35..c44fba2c69 100644 --- a/.github/workflows/build-cmake-pkg.yml +++ b/.github/workflows/build-cmake-pkg.yml @@ -27,30 +27,26 @@ jobs: cmake --install build --prefix "$PREFIX" --config Release export LLAMA_CONFIG="$PREFIX"/lib/cmake/llama/llama-config.cmake - tclsh <<'EOF' - set build(commit) [string trim [exec git rev-parse --short HEAD]] - set build(number) [string trim [exec git rev-list --count HEAD]] + build_commit=$(git rev-parse --short HEAD | xargs) + build_number=$(git rev-list --count HEAD | xargs) - set cmakelists [read [open "CMakeLists.txt" r]] - regexp {set\(LLAMA_VERSION_MAJOR\s+(\d+)\)} $cmakelists -> major - regexp {set\(LLAMA_VERSION_MINOR\s+(\d+)\)} $cmakelists -> minor - regexp {set\(LLAMA_VERSION_PATCH\s+(\d+)\)} $cmakelists -> patch - set build(version) "$major.$minor.$patch" + major=$(grep -oE "set\(LLAMA_VERSION_MAJOR[[:space:]]+[0-9]+" CMakeLists.txt | grep -oE "[0-9]+$") + minor=$(grep -oE "set\(LLAMA_VERSION_MINOR[[:space:]]+[0-9]+" CMakeLists.txt | grep -oE "[0-9]+$") + patch=$(grep -oE "set\(LLAMA_VERSION_PATCH[[:space:]]+[0-9]+" CMakeLists.txt | grep -oE "[0-9]+$") + build_version="$major.$minor.$patch" - set llamaconfig [read [open "$env(LLAMA_CONFIG)" r]] - set checks [list "set\\(LLAMA_VERSION \\s+$build(version)\\)" \ - "set\\(LLAMA_BUILD_COMMIT\\s+$build(commit)\\)" \ - "set\\(LLAMA_BUILD_NUMBER\\s+$build(number)\\)"] + checks=("set\(LLAMA_VERSION[[:space:]]+$build_version\)" + "set\(LLAMA_BUILD_COMMIT[[:space:]]+$build_commit\)" + "set\(LLAMA_BUILD_NUMBER[[:space:]]+$build_number\)") - puts -nonewline "Checking llama-config.cmake version... " - foreach check $checks { - if {![regexp -expanded -- $check $llamaconfig]} { - puts "\"$check\" failed!" + for check in "${checks[@]}"; do + if ! grep -qE "$check" "$LLAMA_CONFIG"; then + echo "Checking llama-config.cmake version... \"$check\" failed!" exit 1 - } - } - puts "success." - EOF + fi + done + + echo "Checking llama-config.cmake version... success." cd examples/simple-cmake-pkg cmake -S . -B build -DCMAKE_PREFIX_PATH="$PREFIX"/lib/cmake From 749f688fcaa4c472ec034b08cb8a907c45cfaa02 Mon Sep 17 00:00:00 2001 From: Xuan-Son Nguyen Date: Fri, 21 Aug 2026 00:36:57 +0200 Subject: [PATCH 02/38] ggml: support ggml_rope_set_offset on opencl, sycl, wgpu, hexagon (#27345) * ggml: support ggml_rope_set_offset on opencl, sycl, wgpu, hexagon * rm inplace optimization --- ggml/src/ggml-hexagon/ggml-hexagon.cpp | 5 +- ggml/src/ggml-hexagon/htp/rope-ops.c | 22 ++-- ggml/src/ggml-opencl/ggml-opencl.cpp | 11 +- ggml/src/ggml-opencl/kernels/rope.cl | 92 +++++++++-------- ggml/src/ggml-sycl/ggml-sycl.cpp | 2 - ggml/src/ggml-sycl/rope.cpp | 106 +++++++++++--------- ggml/src/ggml-webgpu/ggml-webgpu.cpp | 8 +- ggml/src/ggml-webgpu/wgsl-shaders/rope.wgsl | 18 ++-- 8 files changed, 152 insertions(+), 112 deletions(-) diff --git a/ggml/src/ggml-hexagon/ggml-hexagon.cpp b/ggml/src/ggml-hexagon/ggml-hexagon.cpp index b262a73d95..e8a5009b38 100644 --- a/ggml/src/ggml-hexagon/ggml-hexagon.cpp +++ b/ggml/src/ggml-hexagon/ggml-hexagon.cpp @@ -3180,8 +3180,9 @@ static bool ggml_hexagon_supported_argsort(const struct ggml_hexagon_session * s static bool ggml_hexagon_supported_rope(const struct ggml_hexagon_session * sess, const struct ggml_tensor * op) { const int32_t * op_params = &op->op_params[0]; - if (op_params[15] != 0) { - return false; // FIXME: support ggml_rope_set_offset + // ggml_rope_set_offset: HVX kernels need a VLEN-aligned window start (32 f32 elems) + if (op_params[15] % 32 != 0) { + return false; } int mode = op_params[2]; diff --git a/ggml/src/ggml-hexagon/htp/rope-ops.c b/ggml/src/ggml-hexagon/htp/rope-ops.c index 5bc7d74f5e..6c68982493 100644 --- a/ggml/src/ggml-hexagon/htp/rope-ops.c +++ b/ggml/src/ggml-hexagon/htp/rope-ops.c @@ -53,6 +53,7 @@ struct htp_rope_context { int32_t n_dims; + int32_t n_offs; int32_t mode; int32_t n_ctx_orig; int32_t sections[4]; @@ -405,32 +406,40 @@ static inline void hvx_rope_f32_aa(float * restrict dst, const float * restrict static void inline rope_basic_f32(struct htp_rope_context * rctx, uint8_t * restrict dst, uint8_t * restrict src, uint32_t nr, uint32_t ne0, const float * restrict theta_cache) { + const uint32_t n_offs = rctx->n_offs; // VLEN-aligned (enforced by supports_op) #pragma unroll(4) for (uint32_t i = 0; i < nr; i++) { float * d = (float *) (dst + i * rctx->dst_row_size_aligned); float * s = (float *) (src + i * rctx->src0_row_size_aligned); - hvx_rope_f32_aa(d, s, rctx->n_dims, theta_cache); + hvx_rope_f32_aa(d + n_offs, s + n_offs, rctx->n_dims, theta_cache); // fill the remain channels with data from src tensor - if (rctx->n_dims < ne0) { - hvx_copy_f32_uu((uint8_t *)(d + rctx->n_dims), (uint8_t *)(s + rctx->n_dims), ne0 - rctx->n_dims); + if (n_offs > 0) { + hvx_copy_f32_uu((uint8_t *) d, (uint8_t *) s, n_offs); + } + if (n_offs + rctx->n_dims < ne0) { + hvx_copy_f32_uu((uint8_t *)(d + n_offs + rctx->n_dims), (uint8_t *)(s + n_offs + rctx->n_dims), ne0 - n_offs - rctx->n_dims); } } } static void inline rope_neox_f32(struct htp_rope_context * rctx, uint8_t * restrict dst, uint8_t * restrict src, uint32_t nr, uint32_t ne0, const float * restrict theta_cache) { + const uint32_t n_offs = rctx->n_offs; // VLEN-aligned (enforced by supports_op) #pragma unroll(4) for (uint32_t i = 0; i < nr; i++) { float * d = (float *) (dst + i * rctx->dst_row_size_aligned); float * s = (float *) (src + i * rctx->src0_row_size_aligned); - hvx_rope_neox_f32_aa(d, s, rctx->n_dims, theta_cache); + hvx_rope_neox_f32_aa(d + n_offs, s + n_offs, rctx->n_dims, theta_cache); // fill the remain channels with data from src tensor - if (rctx->n_dims < ne0) { - hvx_copy_f32_uu((uint8_t *)(d + rctx->n_dims), (uint8_t *)(s + rctx->n_dims), ne0 - rctx->n_dims); + if (n_offs > 0) { + hvx_copy_f32_uu((uint8_t *) d, (uint8_t *) s, n_offs); + } + if (n_offs + rctx->n_dims < ne0) { + hvx_copy_f32_uu((uint8_t *)(d + n_offs + rctx->n_dims), (uint8_t *)(s + n_offs + rctx->n_dims), ne0 - n_offs - rctx->n_dims); } } } @@ -673,6 +682,7 @@ static int execute_op_rope_f32(struct htp_ops_context * octx) { rctx.n_dims = ((const int32_t *) op_params)[1]; rctx.mode = ((const int32_t *) op_params)[2]; rctx.n_ctx_orig = ((const int32_t *) op_params)[4]; + rctx.n_offs = ((const int32_t *) op_params)[15]; memcpy(&rctx.freq_base, (int32_t *) op_params + 5, sizeof(float)); memcpy(&rctx.freq_scale, (int32_t *) op_params + 6, sizeof(float)); diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 49cd9fd355..26f952a173 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -7434,9 +7434,6 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te case GGML_OP_DIAG_MASK_INF: return op->ne[3] == 1; case GGML_OP_ROPE: { - if (((const int32_t *) op->op_params)[15] != 0) { - return false; // FIXME: support ggml_rope_set_offset - } const int mode = ((const int32_t *) op->op_params)[2]; const bool is_mrope = mode & GGML_ROPE_TYPE_MROPE; const bool is_vision = mode == GGML_ROPE_TYPE_VISION; @@ -23910,6 +23907,7 @@ static void ggml_cl_rope(ggml_backend_t backend, const ggml_tensor * src0, const const int n_dims = ((int *) dst->op_params)[1]; const int mode = ((int *) dst->op_params)[2]; const int n_ctx_orig = ((int32_t *) dst->op_params)[4]; + const int n_offs = ((int32_t *) dst->op_params)[15]; float freq_base; float freq_scale; @@ -23938,6 +23936,7 @@ static void ggml_cl_rope(ggml_backend_t backend, const ggml_tensor * src0, const if (is_vision) { GGML_ASSERT(n_dims == ne00/2); + GGML_ASSERT(n_offs == 0); // offset not supported for vision, as the rotated pairs span the whole row } cl_kernel kernel; @@ -24029,6 +24028,12 @@ static void ggml_cl_rope(ggml_backend_t backend, const ggml_tensor * src0, const if (is_mrope && !is_vision) { CL_CHECK(clSetKernelArg(kernel, 34, sizeof(int), &is_imrope)); } + // norm and neox have n_offs after beta_slow, mrope has it after is_imrope + if (!is_mrope && !is_vision) { + CL_CHECK(clSetKernelArg(kernel, 33, sizeof(int), &n_offs)); + } else if (is_mrope && !is_vision) { + CL_CHECK(clSetKernelArg(kernel, 35, sizeof(int), &n_offs)); + } size_t global_work_size[] = {(size_t)ne01*nth, (size_t)ne02, (size_t)ne03}; size_t local_work_size[] = {(size_t)nth, 1, 1}; diff --git a/ggml/src/ggml-opencl/kernels/rope.cl b/ggml/src/ggml-opencl/kernels/rope.cl index 82f4cd8740..27fdbbbc4f 100644 --- a/ggml/src/ggml-opencl/kernels/rope.cl +++ b/ggml/src/ggml-opencl/kernels/rope.cl @@ -75,7 +75,8 @@ kernel void kernel_rope_norm_f32( float ext_factor, float attn_factor, float beta_fast, - float beta_slow + float beta_slow, + int n_offs ) { src0 = (global void*)((global char*)src0 + offset0); src1 = (global int*)((global char*)src1 + offset1); @@ -94,14 +95,15 @@ kernel void kernel_rope_norm_f32( float inv_ndims = -1.f/n_dims; for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) { - if (i0 < n_dims) { - int ic = i0/2; + if (i0 >= n_offs && i0 < n_offs + n_dims) { + int iw = i0 - n_offs; // relative idx + int ic = iw/2; - float theta = theta_base * pow(freq_base, inv_ndims*i0); + float theta = theta_base * pow(freq_base, inv_ndims*iw); float freq_factor = src2 != src0 ? src2[ic] : 1.0f; - float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor); + float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor); global float * src = (global float *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00); global float * dst_data = (global float *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0); @@ -154,7 +156,8 @@ kernel void kernel_rope_norm_f16( float ext_factor, float attn_factor, float beta_fast, - float beta_slow + float beta_slow, + int n_offs ) { src0 = (global void*)((global char*)src0 + offset0); src1 = (global int*)((global char*)src1 + offset1); @@ -173,14 +176,15 @@ kernel void kernel_rope_norm_f16( float inv_ndims = -1.f/n_dims; for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) { - if (i0 < n_dims) { - int ic = i0/2; + if (i0 >= n_offs && i0 < n_offs + n_dims) { + int iw = i0 - n_offs; // relative idx + int ic = iw/2; - float theta = theta_base * pow(freq_base, inv_ndims*i0); + float theta = theta_base * pow(freq_base, inv_ndims*iw); float freq_factor = src2 != src0 ? src2[ic] : 1.0f; - float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor); + float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor); global half * src = (global half *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00); global half * dst_data = (global half *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0); @@ -233,7 +237,8 @@ kernel void kernel_rope_neox_f32( float ext_factor, float attn_factor, float beta_fast, - float beta_slow + float beta_slow, + int n_offs ) { src0 = (global void*)((global char*)src0 + offset0); src1 = (global int*)((global char*)src1 + offset1); @@ -252,17 +257,18 @@ kernel void kernel_rope_neox_f32( float inv_ndims = -1.f/n_dims; for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) { - if (i0 < n_dims) { - int ic = i0/2; + if (i0 >= n_offs && i0 < n_offs + n_dims) { + int iw = i0 - n_offs; // relative idx + int ic = iw/2; - const float theta = theta_base * pow(freq_base, inv_ndims*i0); + const float theta = theta_base * pow(freq_base, inv_ndims*iw); const float freq_factor = src2 != src0 ? src2[ic] : 1.0f; - float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor); + float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor); - global float * src = (global float *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + ic*nb00); - global float * dst_data = (global float *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + ic*nb0); + global float * src = (global float *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + (n_offs + ic)*nb00); + global float * dst_data = (global float *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + (n_offs + ic)*nb0); const float x0 = src[0]; const float x1 = src[n_dims/2]; @@ -312,7 +318,8 @@ kernel void kernel_rope_neox_f16( float ext_factor, float attn_factor, float beta_fast, - float beta_slow + float beta_slow, + int n_offs ) { src0 = (global void*)((global char*)src0 + offset0); src1 = (global int*)((global char*)src1 + offset1); @@ -331,17 +338,18 @@ kernel void kernel_rope_neox_f16( float inv_ndims = -1.f/n_dims; for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) { - if (i0 < n_dims) { - int ic = i0/2; + if (i0 >= n_offs && i0 < n_offs + n_dims) { + int iw = i0 - n_offs; // relative idx + int ic = iw/2; - const float theta = theta_base * pow(freq_base, inv_ndims*i0); + const float theta = theta_base * pow(freq_base, inv_ndims*iw); const float freq_factor = src2 != src0 ? src2[ic] : 1.0f; - float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor); + float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor); - global half * src = (global half *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + ic*nb00); - global half * dst_data = (global half *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + ic*nb0); + global half * src = (global half *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + (n_offs + ic)*nb00); + global half * dst_data = (global half *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + (n_offs + ic)*nb0); const float x0 = src[0]; const float x1 = src[n_dims/2]; @@ -393,7 +401,8 @@ kernel void kernel_rope_multi_f32( float beta_fast, float beta_slow, int4 sections, - int is_imrope + int is_imrope, + int n_offs ) { src0 = (global void*)((global char*)src0 + offset0); src1 = (global int*)((global char*)src1 + offset1); @@ -414,10 +423,11 @@ kernel void kernel_rope_multi_f32( float inv_ndims = -1.f/n_dims; for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) { - if (i0 < n_dims) { - int ic = i0/2; + if (i0 >= n_offs && i0 < n_offs + n_dims) { + int iw = i0 - n_offs; // relative idx + int ic = iw/2; - const int sector = (i0 / 2) % sect_dims; + const int sector = ic % sect_dims; float theta_base = 0.0f; if (is_imrope) { @@ -445,14 +455,14 @@ kernel void kernel_rope_multi_f32( } } - const float theta = theta_base * pow(freq_base, inv_ndims*i0); + const float theta = theta_base * pow(freq_base, inv_ndims*iw); const float freq_factor = src2 != src0 ? src2[ic] : 1.0f; - float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor); + float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor); - global float * src = (global float *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + ic*nb00); - global float * dst_data = (global float *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + ic*nb0); + global float * src = (global float *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + (n_offs + ic)*nb00); + global float * dst_data = (global float *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + (n_offs + ic)*nb0); const float x0 = src[0]; const float x1 = src[n_dims/2]; @@ -504,7 +514,8 @@ kernel void kernel_rope_multi_f16( float beta_fast, float beta_slow, int4 sections, - int is_imrope + int is_imrope, + int n_offs ) { src0 = (global void*)((global char*)src0 + offset0); src1 = (global int*)((global char*)src1 + offset1); @@ -525,10 +536,11 @@ kernel void kernel_rope_multi_f16( float inv_ndims = -1.f/n_dims; for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) { - if (i0 < n_dims) { - int ic = i0/2; + if (i0 >= n_offs && i0 < n_offs + n_dims) { + int iw = i0 - n_offs; // relative idx + int ic = iw/2; - const int sector = (i0 / 2) % sect_dims; + const int sector = ic % sect_dims; float theta_base = 0.0f; if (is_imrope) { @@ -556,14 +568,14 @@ kernel void kernel_rope_multi_f16( } } - const float theta = theta_base * pow(freq_base, inv_ndims*i0); + const float theta = theta_base * pow(freq_base, inv_ndims*iw); const float freq_factor = src2 != src0 ? src2[ic] : 1.0f; - float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor); + float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor); - global half * src = (global half *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + ic*nb00); - global half * dst_data = (global half *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + ic*nb0); + global half * src = (global half *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + (n_offs + ic)*nb00); + global half * dst_data = (global half *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + (n_offs + ic)*nb0); const float x0 = src[0]; const float x1 = src[n_dims/2]; diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index c7434a6bdb..57aae9011d 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -6242,8 +6242,6 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons } case GGML_OP_ROPE: case GGML_OP_ROPE_BACK: - // FIXME: support ggml_rope_set_offset - return ((const int32_t *) op->op_params)[15] == 0; case GGML_OP_IM2COL: case GGML_OP_IM2COL_3D: case GGML_OP_UPSCALE: diff --git a/ggml/src/ggml-sycl/rope.cpp b/ggml/src/ggml-sycl/rope.cpp index 9d83a1e9fa..b6d22559d1 100644 --- a/ggml/src/ggml-sycl/rope.cpp +++ b/ggml/src/ggml-sycl/rope.cpp @@ -41,7 +41,7 @@ template static void rope_norm(const T *x, D *dst, const int ne00, const int ne01, const int ne02, const int s01, const int s02, const int s03, const int s1, const int s2, const int s3, - const int n_dims, const int32_t *pos, + const int n_dims, const int n_offs, 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, @@ -78,19 +78,21 @@ static void rope_norm(const T *x, D *dst, const int ne00, const int ne01, ggml_sycl_memcpy_1<4>(dst + idst, &v); } }; - if (i0 >= n_dims) { + if (i0 < n_offs || i0 >= n_offs + n_dims) { store_coaelsced(x[ix + 0], x[ix + 1]); return; } - const float theta_base = pos[i2] * dpct::pow(theta_scale, i0 / 2.0f); + const int iw = i0 - n_offs; // relative idx - const float freq_factor = has_ff ? freq_factors[i0 / 2] : 1.0f; + const float theta_base = pos[i2] * dpct::pow(theta_scale, iw / 2.0f); + + const float freq_factor = has_ff ? freq_factors[iw / 2] : 1.0f; float cos_theta; float sin_theta; - rope_yarn(theta_base / freq_factor, freq_scale, corr_dims, i0, + rope_yarn(theta_base / freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor, cos_theta, sin_theta); const float x0 = x[ix + 0]; @@ -104,7 +106,7 @@ template static void rope_neox(const T *x, D *dst, const int ne00, const int ne01, const int ne02, const int s01, const int s02, const int s03, const int s1, const int s2, const int s3, - const int n_dims, const int32_t *pos, + const int n_dims, const int n_offs, 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, @@ -132,35 +134,38 @@ static void rope_neox(const T *x, D *dst, const int ne00, const int ne01, idst += row_indices[i2] * set_rows_stride; } - if (i0 >= n_dims) { + if (i0 < n_offs || i0 >= n_offs + n_dims) { dst[idst + i0 / 2 + 0] = ggml_sycl_cast(x[ix + i0 / 2 + 0]); dst[idst + i0 / 2 + 1] = ggml_sycl_cast(x[ix + i0 / 2 + 1]); return; } - const float theta_base = pos[i2] * dpct::pow(theta_scale, i0 / 2.0f); + const int iw = i0 - n_offs; // relative idx - const float freq_factor = has_ff ? freq_factors[i0 / 2] : 1.0f; + const float theta_base = pos[i2] * dpct::pow(theta_scale, iw / 2.0f); + + const float freq_factor = has_ff ? freq_factors[iw / 2] : 1.0f; float cos_theta; float sin_theta; - rope_yarn(theta_base / freq_factor, freq_scale, corr_dims, i0, + rope_yarn(theta_base / freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor, cos_theta, sin_theta); - const float x0 = x[ix + 0]; - const float x1 = x[ix + n_dims / 2]; + // idst/ix point at channel i0/2; the first channel of the rotated pair is n_offs + iw/2 = i0/2 + n_offs/2 + const float x0 = x[ix + n_offs / 2 + 0]; + const float x1 = x[ix + n_offs / 2 + n_dims / 2]; - dst[idst + 0] = ggml_sycl_cast(x0 * cos_theta - x1 * sin_theta); - dst[idst + n_dims / 2] = ggml_sycl_cast(x0 * sin_theta + x1 * cos_theta); + dst[idst + n_offs / 2 + 0] = ggml_sycl_cast(x0 * cos_theta - x1 * sin_theta); + dst[idst + n_offs / 2 + n_dims / 2] = ggml_sycl_cast(x0 * sin_theta + x1 * cos_theta); } template static void rope_multi(const T *x, T *dst, const int ne00, const int ne01, const int ne02, const int s01, const int s02, const int s03, const int s1, const int s2, const int s3, - const int n_dims, const int32_t *pos, + const int n_dims, const int n_offs, 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, @@ -183,54 +188,57 @@ static void rope_multi(const T *x, T *dst, const int ne00, const int ne01, int idst = i0 / 2 + i1 * s1 + i2 * s2 + i3 * s3; const int ix = i0 / 2 + i1 * s01 + i2 * s02 + i3 * s03; - if (i0 >= n_dims) { + if (i0 < n_offs || i0 >= n_offs + n_dims) { dst[idst + i0 / 2 + 0] = x[ix + i0 / 2 + 0]; dst[idst + i0 / 2 + 1] = x[ix + i0 / 2 + 1]; return; } + const int iw = i0 - n_offs; // relative idx + const int sect_dims = sections.v[0] + sections.v[1] + sections.v[2] + sections.v[3]; const int sec_w = sections.v[1] + sections.v[0]; - const int sector = (i0 / 2) % sect_dims; + const int sector = (iw / 2) % sect_dims; float theta_base = 0.0; if (is_imrope) { if (sector % 3 == 1 && sector < 3 * sections.v[1]) { // h - theta_base = pos[i2 + ne02 * 1] * dpct::pow(theta_scale, i0 / 2.0f); + theta_base = pos[i2 + ne02 * 1] * dpct::pow(theta_scale, iw / 2.0f); } else if (sector % 3 == 2 && sector < 3 * sections.v[2]) { // w - theta_base = pos[i2 + ne02 * 2] * dpct::pow(theta_scale, i0 / 2.0f); + theta_base = pos[i2 + ne02 * 2] * dpct::pow(theta_scale, iw / 2.0f); } else if (sector % 3 == 0 && sector < 3 * sections.v[0]) { // t - theta_base = pos[i2] * dpct::pow(theta_scale, i0 / 2.0f); + theta_base = pos[i2] * dpct::pow(theta_scale, iw / 2.0f); } else { - theta_base = pos[i2 + ne02 * 3] * dpct::pow(theta_scale, i0 / 2.0f); + theta_base = pos[i2 + ne02 * 3] * dpct::pow(theta_scale, iw / 2.0f); } } else { if (sector < sections.v[0]) { - theta_base = pos[i2] * dpct::pow(theta_scale, i0 / 2.0f); + theta_base = pos[i2] * dpct::pow(theta_scale, iw / 2.0f); } else if (sector >= sections.v[0] && sector < sec_w) { - theta_base = pos[i2 + ne02 * 1] * dpct::pow(theta_scale, i0 / 2.0f); + theta_base = pos[i2 + ne02 * 1] * dpct::pow(theta_scale, iw / 2.0f); } else if (sector >= sec_w && sector < sec_w + sections.v[2]) { - theta_base = pos[i2 + ne02 * 2] * dpct::pow(theta_scale, i0 / 2.0f); + theta_base = pos[i2 + ne02 * 2] * dpct::pow(theta_scale, iw / 2.0f); } else if (sector >= sec_w + sections.v[2]) { - theta_base = pos[i2 + ne02 * 3] * dpct::pow(theta_scale, i0 / 2.0f); + theta_base = pos[i2 + ne02 * 3] * dpct::pow(theta_scale, iw / 2.0f); } } - const float freq_factor = has_ff ? freq_factors[i0 / 2] : 1.0f; + const float freq_factor = has_ff ? freq_factors[iw / 2] : 1.0f; float cos_theta; float sin_theta; - rope_yarn(theta_base / freq_factor, freq_scale, corr_dims, i0, + rope_yarn(theta_base / freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor, cos_theta, sin_theta); - const float x0 = x[ix + 0]; - const float x1 = x[ix + n_dims / 2]; + // idst/ix point at channel i0/2; the first channel of the rotated pair is n_offs + iw/2 = i0/2 + n_offs/2 + const float x0 = x[ix + n_offs / 2 + 0]; + const float x1 = x[ix + n_offs / 2 + 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 + n_offs / 2 + 0] = x0 * cos_theta - x1 * sin_theta; + dst[idst + n_offs / 2 + n_dims / 2] = x0 * sin_theta + x1 * cos_theta; } template @@ -293,7 +301,7 @@ static void rope_norm_sycl(const T *x, D *dst, const int ne00, const int ne01, const int ne02, const int s01, const int s02, const int s03, const int s1, const int s2, const int s3, const int n_dims, - const int nr, const int32_t *pos, const float freq_scale, + const int n_offs, 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, @@ -313,7 +321,7 @@ rope_norm_sycl(const T *x, D *dst, const int ne00, const int ne01, GGML_UNUSED(item_ct1); rope_norm( x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, - pos, freq_scale, ext_factor, attn_factor, corr_dims, + n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride); }); } else { @@ -323,7 +331,7 @@ rope_norm_sycl(const T *x, D *dst, const int ne00, const int ne01, GGML_UNUSED(item_ct1); rope_norm( x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, - pos, freq_scale, ext_factor, attn_factor, corr_dims, + n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride); }); } @@ -334,7 +342,7 @@ static void rope_neox_sycl(const T *x, D *dst, const int ne00, const int ne01, const int ne02, const int s01, const int s02, const int s03, const int s1, const int s2, const int s3, const int n_dims, - const int nr, const int32_t *pos, const float freq_scale, + const int n_offs, 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, @@ -354,7 +362,7 @@ rope_neox_sycl(const T *x, D *dst, const int ne00, const int ne01, GGML_UNUSED(item_ct1); rope_neox( x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, - pos, freq_scale, ext_factor, attn_factor, corr_dims, + n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride); }); } else { @@ -364,7 +372,7 @@ rope_neox_sycl(const T *x, D *dst, const int ne00, const int ne01, GGML_UNUSED(item_ct1); rope_neox( x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, - pos, freq_scale, ext_factor, attn_factor, corr_dims, + n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale, freq_factors, row_indices, set_rows_stride); }); } @@ -375,7 +383,7 @@ static void rope_multi_sycl(const T *x, T *dst, const int ne00, const int ne01, const int ne02, const int s01, const int s02, const int s03, const int s1, const int s2, const int s3, const int n_dims, - const int nr, const int32_t *pos, const float freq_scale, + const int n_offs, 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 mrope_sections sections, @@ -395,7 +403,7 @@ rope_multi_sycl(const T *x, T *dst, const int ne00, const int ne01, GGML_UNUSED(item_ct1); rope_multi( x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, - pos, freq_scale, ext_factor, attn_factor, corr_dims, + n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale, freq_factors, sections, is_imrope); }); } else { @@ -405,7 +413,7 @@ rope_multi_sycl(const T *x, T *dst, const int ne00, const int ne01, GGML_UNUSED(item_ct1); rope_multi( x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, - pos, freq_scale, ext_factor, attn_factor, corr_dims, + n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims, theta_scale, freq_factors, sections, is_imrope); }); } @@ -497,6 +505,7 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst, const int n_dims = ((int32_t *)dst->op_params)[1]; const int mode = ((int32_t *)dst->op_params)[2]; const int n_ctx_orig = ((int32_t *)dst->op_params)[4]; + const int n_offs = ((int32_t *)dst->op_params)[15]; mrope_sections sections; float freq_base; @@ -526,6 +535,7 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst, if (is_vision) { GGML_ASSERT(n_dims == ne00 / 2); + GGML_ASSERT(n_offs == 0); // offset not supported for vision, as the rotated pairs span the whole row } const int32_t *pos = (const int32_t *)src1_d; @@ -545,19 +555,19 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst, if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F32) { rope_neox_sycl( (const float *)src0_d, (float *)dst_d, ne00, ne01, ne02, s01, - s02, s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base, + s02, s03, s1, s2, s3, n_dims, n_offs, 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_sycl( (const float *)src0_d, (sycl::half *)dst_d, ne00, ne01, ne02, - s01, s02, s03, s1, s2, s3, n_dims, nr, pos, freq_scale, + s01, s02, s03, s1, s2, s3, n_dims, n_offs, 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_sycl( (const sycl::half *)src0_d, (sycl::half *)dst_d, ne00, ne01, - ne02, s01, s02, s03, s1, s2, s3, n_dims, nr, pos, freq_scale, + ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base, ext_factor, attn_factor, corr_dims, freq_factors, row_indices, set_rows_stride, stream); } else { @@ -568,13 +578,13 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst, if (src0->type == GGML_TYPE_F32) { rope_multi_sycl((const float *)src0_d, (float *)dst_d, ne00, ne01, ne02, s01, s02, s03, s1, s2, - s3, n_dims, nr, pos, freq_scale, freq_base, + s3, n_dims, n_offs, nr, pos, freq_scale, freq_base, ext_factor, attn_factor, corr_dims, freq_factors, sections, is_imrope, stream); } else if (src0->type == GGML_TYPE_F16) { rope_multi_sycl( (const sycl::half *)src0_d, (sycl::half *)dst_d, ne00, ne01, - ne02, s01, s02, s03, s1, s2, s3, n_dims, nr, pos, freq_scale, + ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base, ext_factor, attn_factor, corr_dims, freq_factors, sections, is_imrope, stream); } else { @@ -602,19 +612,19 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst, if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F32) { rope_norm_sycl( (const float *)src0_d, (float *)dst_d, ne00, ne01, ne02, s01, - s02, s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base, + s02, s03, s1, s2, s3, n_dims, n_offs, 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_sycl( (const float *)src0_d, (sycl::half *)dst_d, ne00, ne01, ne02, - s01, s02, s03, s1, s2, s3, n_dims, nr, pos, freq_scale, + s01, s02, s03, s1, s2, s3, n_dims, n_offs, 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_sycl( (const sycl::half *)src0_d, (sycl::half *)dst_d, ne00, ne01, - ne02, s01, s02, s03, s1, s2, s3, n_dims, nr, pos, freq_scale, + ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base, ext_factor, attn_factor, corr_dims, freq_factors, row_indices, set_rows_stride, stream); } else { diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp index 4367f9a610..2434848a55 100644 --- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp +++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp @@ -2714,6 +2714,7 @@ static webgpu_encoded_op ggml_webgpu_rope(webgpu_context & ctx, const int n_dims = ((int32_t *) dst->op_params)[1]; const int mode = ((int32_t *) dst->op_params)[2]; const int n_ctx_orig = ((int32_t *) dst->op_params)[4]; + const int n_offs = ((int32_t *) dst->op_params)[15]; float freq_base; float freq_scale; @@ -2762,7 +2763,8 @@ static webgpu_encoded_op ggml_webgpu_rope(webgpu_context & ctx, (uint32_t) sections[0], (uint32_t) sections[1], (uint32_t) sections[2], - (uint32_t) sections[3] + (uint32_t) sections[3], + (uint32_t) n_offs }; std::vector entries = { ggml_webgpu_make_tensor_bind_group_entry(ctx, 0, src0), @@ -4472,9 +4474,7 @@ static bool ggml_backend_webgpu_device_supports_op(ggml_backend_dev_t dev, const supports_op = (op->type == GGML_TYPE_F32 && src0->type == GGML_TYPE_F32) && ggml_is_contiguous_rows(src0); break; case GGML_OP_ROPE: - // FIXME: support ggml_rope_set_offset - supports_op = - (op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16) && ((const int32_t *) op->op_params)[15] == 0; + supports_op = op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16; break; case GGML_OP_GLU: switch (ggml_get_glu_op(op)) { diff --git a/ggml/src/ggml-webgpu/wgsl-shaders/rope.wgsl b/ggml/src/ggml-webgpu/wgsl-shaders/rope.wgsl index 1c874e1424..6ff53088c4 100644 --- a/ggml/src/ggml-webgpu/wgsl-shaders/rope.wgsl +++ b/ggml/src/ggml-webgpu/wgsl-shaders/rope.wgsl @@ -38,7 +38,8 @@ struct Params { sections0: u32, sections1: u32, sections2: u32, - sections3: u32 + sections3: u32, + n_offs: u32 }; @group(0) @binding(0) @@ -126,7 +127,8 @@ fn rope_yarn(theta_extrap: f32, i: u32) -> vec2 { fn pair_base(i0: u32, div_2: bool) -> u32 { if (div_2) { - return i0 / 2; + // first channel of the rotated pair: n_offs + (i0 - n_offs)/2 + return i0 / 2 + params.n_offs / 2; } else { return i0; } @@ -165,20 +167,22 @@ fn main(@builtin(global_invocation_id) gid: vec3) { let i_src_row = params.offset_src0 + i3 * params.stride_src03 + i2 * params.stride_src02 + i1 * params.stride_src01; let i_dst_row = params.offset_dst + i3 * params.stride_dst3 + i2 * params.stride_dst2 + i1 * params.stride_dst1; - if (i0 >= params.n_dims && !is_vision) { + if ((i0 < params.n_offs || i0 >= params.n_offs + params.n_dims) && !is_vision) { let i_src = i_src_row + i0; let i_dst = i_dst_row + i0; rotate(i_dst, i_dst + 1, f32(src0[i_src]), f32(src0[i_src + 1])); return; } + let iw = i0 - params.n_offs; // relative idx + var theta_base_mult: u32 = 0; - var theta_scale_pwr: u32 = i0 / 2; + var theta_scale_pwr: u32 = iw / 2; if (is_mrope) { let sect_dims = params.sections0 + params.sections1 + params.sections2 + params.sections3; let sec_w = params.sections1 + params.sections0; let sec_e = params.sections2 + sec_w; - let sector = (i0 / 2) % sect_dims; + let sector = (iw / 2) % sect_dims; if (is_imrope) { if (sector % 3 == 1 && sector < 3 * params.sections1) { theta_base_mult = 1; @@ -203,7 +207,7 @@ fn main(@builtin(global_invocation_id) gid: vec3) { } else if (sector >= sec_e) { if (is_vision) { theta_scale_pwr = sector - sec_e; - theta_scale_pwr = (i0 / 2) % sec_e; + theta_scale_pwr = (iw / 2) % sec_e; } theta_base_mult = 3; } else if (is_vision) { @@ -212,7 +216,7 @@ fn main(@builtin(global_invocation_id) gid: vec3) { } } let theta_base = f32(src1[params.offset_src1 + i2 + params.ne2 * theta_base_mult]) * pow(params.theta_scale, f32(theta_scale_pwr)); - let thetas = rope_yarn(theta_base/freq_factor(i0), i0); + let thetas = rope_yarn(theta_base/freq_factor(iw), iw); let i_src = i_src_row + pair_base(i0, is_neox || is_mrope || is_vision); let i_dst = i_dst_row + pair_base(i0, is_neox || is_mrope || is_vision); From a298422da78eb75e440a7de0ca408af64d323d93 Mon Sep 17 00:00:00 2001 From: vk <89937361+itsvedantkumar@users.noreply.github.com> Date: Fri, 21 Aug 2026 10:06:59 +0530 Subject: [PATCH 03/38] docs: fix typos in ET.md (#27457) --- docs/backend/ET.md | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/docs/backend/ET.md b/docs/backend/ET.md index 8d9ba12c82..8ebc15fb7d 100644 --- a/docs/backend/ET.md +++ b/docs/backend/ET.md @@ -116,7 +116,7 @@ in inline assembler. Most kernels are very naive with lots of low hanging fruits left: > [!IMPORTANT] -> Several assembly instructions emmited by the compiler are not implemented +> Several assembly instructions emitted by the compiler are not implemented > in hardware and software emulation in firmware is not ready yet. > Eventually firmware will transparently trap unimplemented instructions > and will emulate them inside exception handler. Until then, kernel @@ -138,12 +138,12 @@ Most kernels are very naive with lots of low hanging fruits left: > kernel build process. Feel free to take ideas/code from there or try linking > it in. -Before commiting any changes to operations and/or kernels, don't forget +Before committing any changes to operations and/or kernels, don't forget to update supported ops reports (instructions at `docs/ops.md`). When logging is enabled (e.g. by setting `--log-file` cli param), each compute kernel run outputs a line with -pipe-delimited key-value pairs containing kernel level performance infomation. +pipe-delimited key-value pairs containing kernel level performance information. Line is prefixed with `ET_PERF`: ``` @@ -160,7 +160,7 @@ to `GGML_ET_PROFILE/et_runtime_trace.json` and `GGML_ET_PROFILE/kernel_map` on e ### Uberkernel -The in-knernel implementaiton of device dispatch/kernel fusion. The ET SDK has a non-trivial op-to-op gap. `Uberkernel` (name taken from the original Esperanto AI's compiler) +The in-kernel implementation of device dispatch/kernel fusion. The ET SDK has a non-trivial op-to-op gap. `Uberkernel` (name taken from the original Esperanto AI's compiler) dispatches multiple already existing kernel implementations with device side synchronization. Due to the processor's design, there is no natural memory visibility horizon between sub-kernel invocations. This makes uberkernel much more difficult to develop and debug. Currently Uberkerel is hidden begind the `GGML_ET_UBERKERNEL` environment variable and is disabled by default. Setting it to 1 enables it and provides significant performance improvements but is only From b2e5e9b28b2484fbf94b543432ece638996a8b97 Mon Sep 17 00:00:00 2001 From: Chris Danis Date: Fri, 21 Aug 2026 01:13:58 -0400 Subject: [PATCH 04/38] TP: enable tensor split for LFM2/LFM2MOE (#26993) Assisted-by: deepseek-v4-flash --- src/llama-arch.cpp | 2 -- src/llama-model.cpp | 4 ++++ 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/src/llama-arch.cpp b/src/llama-arch.cpp index 4089544016..c9b504c333 100644 --- a/src/llama-arch.cpp +++ b/src/llama-arch.cpp @@ -1062,8 +1062,6 @@ bool llm_arch_supports_sm_tensor(const llm_arch & arch) { case LLM_ARCH_NEMOTRON_H: case LLM_ARCH_NEMOTRON_H_MOE: case LLM_ARCH_GRANITE_HYBRID: - case LLM_ARCH_LFM2: - case LLM_ARCH_LFM2MOE: case LLM_ARCH_MINIMAX_01: case LLM_ARCH_MINIMAX_M2: case LLM_ARCH_MINIMAX_M3: diff --git a/src/llama-model.cpp b/src/llama-model.cpp index 3759c86259..d7874e0a92 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -487,6 +487,10 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_1, "ssm_out.weight"); } if (std::regex_match(tensor_name, pattern_r_cache) || std::regex_match(tensor_name, pattern_s_cache)) { + if (ud->model->arch == LLM_ARCH_LFM2 || ud->model->arch == LLM_ARCH_LFM2MOE) { + // the LFM2 shortconv block runs fully mirrored, so its conv state must be mirrored too + return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_MIRRORED, ""); + } return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_0, "ssm_out.weight"); } if (std::regex_match(tensor_name, pattern_ssm_conv1d)) { From 9e96cf77ffd4ebd05bad82932906ec3f59ed54ce Mon Sep 17 00:00:00 2001 From: Neo Zhang Date: Fri, 21 Aug 2026 13:14:54 +0800 Subject: [PATCH 05/38] sycl : fix load model with mlock issue (#27250) --- ggml/src/ggml-sycl/ggml-sycl.cpp | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 57aae9011d..8372428795 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -1517,8 +1517,13 @@ static ggml_backend_buffer_t ggml_backend_sycl_host_buffer_type_alloc_buffer(ggm } static size_t ggml_backend_sycl_host_buffer_type_get_max_size(ggml_backend_buffer_type_t buft) { - ggml_backend_sycl_device_context * dev_ctx = (ggml_backend_sycl_device_context *) buft->device->context; - return dpct::dev_mgr::instance().get_device(dev_ctx->device).get_max_mem_alloc_size(); + + if (g_ggml_sycl_enable_host_pinned_mem) { + ggml_backend_sycl_device_context * dev_ctx = (ggml_backend_sycl_device_context *) buft->device->context; + return dpct::dev_mgr::instance().get_device(dev_ctx->device).get_max_mem_alloc_size(); + } else { + return SIZE_MAX; + } } ggml_backend_buffer_type_t ggml_backend_sycl_host_buffer_type() { From 6602dd338941d1fd562e6ca3934acc5602c15a9d Mon Sep 17 00:00:00 2001 From: Ian Faust Date: Fri, 21 Aug 2026 07:15:40 +0200 Subject: [PATCH 06/38] sycl: fix multiple warnings in compiling sycl backend (#26713) * Update norm.cpp * Update helper.hpp * Update im2col.cpp * Update fattn-mkl.cpp * Update element_wise.cpp * Update fattn-mkl.cpp * Update set_rows.cpp * Update element_wise.cpp * Update ggml-sycl.cpp * Update ggml-sycl.cpp * Update ggml-sycl.cpp * Update ggml-sycl.cpp * Update ggml-sycl.cpp * Update norm.cpp * Update CMakeLists.txt * Update CMakeLists.txt * Update CMakeLists.txt * Update ggml-sycl.cpp --- ggml/src/ggml-cpu/CMakeLists.txt | 9 +++++---- ggml/src/ggml-sycl/dpct/helper.hpp | 2 +- ggml/src/ggml-sycl/element_wise.cpp | 8 ++++---- ggml/src/ggml-sycl/fattn-mkl.cpp | 9 ++++----- ggml/src/ggml-sycl/ggml-sycl.cpp | 20 +++++++++++--------- ggml/src/ggml-sycl/im2col.cpp | 4 ++-- ggml/src/ggml-sycl/norm.cpp | 8 -------- ggml/src/ggml-sycl/set_rows.cpp | 2 +- 8 files changed, 28 insertions(+), 34 deletions(-) diff --git a/ggml/src/ggml-cpu/CMakeLists.txt b/ggml/src/ggml-cpu/CMakeLists.txt index a6cc49586b..32e1e7aa19 100644 --- a/ggml/src/ggml-cpu/CMakeLists.txt +++ b/ggml/src/ggml-cpu/CMakeLists.txt @@ -737,8 +737,9 @@ function(ggml_add_cpu_backend_variant_impl tag_name) set_target_properties(${GGML_CPU_NAME} PROPERTIES COMPILE_FLAGS "-msimd128") endif() - if (CMAKE_CXX_COMPILER_ID STREQUAL "IntelLLVM") - # The compiler automatically enables "-ffast-math" which can cause NaNs in tests due to "-fassociative-math" - target_compile_options(${GGML_CPU_NAME} PRIVATE "-fno-associative-math") - endif() + if (CMAKE_C_COMPILER_ID STREQUAL "IntelLLVM" OR CMAKE_CXX_COMPILER_ID STREQUAL "IntelLLVM") + # The compiler automatically enables "-ffast-math" which can cause NaNs in tests due to "-fassociative-math" + target_compile_options(${GGML_CPU_NAME} PRIVATE "$<$,$>:$<$:/clang:>-fno-associative-math>") + endif() + endfunction() diff --git a/ggml/src/ggml-sycl/dpct/helper.hpp b/ggml/src/ggml-sycl/dpct/helper.hpp index 664b8e9697..85af4cab68 100644 --- a/ggml/src/ggml-sycl/dpct/helper.hpp +++ b/ggml/src/ggml-sycl/dpct/helper.hpp @@ -62,7 +62,7 @@ #define DPCT_UNUSED(x) (void)(x) -inline void _abort(const char * str) { +[[noreturn]] inline void _abort(const char * str) { std::cerr << str << std::endl; std::abort(); } diff --git a/ggml/src/ggml-sycl/element_wise.cpp b/ggml/src/ggml-sycl/element_wise.cpp index 8619ed6f4b..95914873e5 100644 --- a/ggml/src/ggml-sycl/element_wise.cpp +++ b/ggml/src/ggml-sycl/element_wise.cpp @@ -10,7 +10,7 @@ (ITEM.get_local_range(IDX) * ITEM.get_group(IDX) + ITEM.get_local_id(IDX)) static void acc_f32(const char * x, const char * y, float * dst, const int64_t ne, - const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, + const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, @@ -455,7 +455,7 @@ static void unary_mul_sycl(const T * x, const T * g, T * dst, const int64_t k, c namespace ggml_sycl_detail { static void acc_f32_sycl(const char *x, const char *y, float *dst, const int64_t n_elements, - const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, + const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, @@ -466,7 +466,7 @@ static void acc_f32_sycl(const char *x, const char *y, float *dst, sycl::range<3>(1, 1, SYCL_ACC_BLOCK_SIZE)), [=](sycl::nd_item<3> /*item_ct1*/) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { acc_f32(x, y, dst, n_elements, - ne0, ne1, ne2, ne3, + ne0, ne1, ne2, nb00, nb01, nb02, nb03, ne10, ne11, ne12, ne13, nb10, nb11, nb12, nb13, @@ -970,7 +970,7 @@ static inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, ggml_tensor const int64_t offset = (int64_t) ((const int32_t *) dst->op_params)[3] / (int64_t) sizeof(float); ggml_sycl_detail::acc_f32_sycl(src0_d, src1_d, dst_d, ggml_nelements(dst), - dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3], + dst->ne[0], dst->ne[1], dst->ne[2], src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3], src1->ne[0], src1->ne[1], src1->ne[2], src1->ne[3], src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3], diff --git a/ggml/src/ggml-sycl/fattn-mkl.cpp b/ggml/src/ggml-sycl/fattn-mkl.cpp index fc22b7bdb8..2d164a0840 100644 --- a/ggml/src/ggml-sycl/fattn-mkl.cpp +++ b/ggml/src/ggml-sycl/fattn-mkl.cpp @@ -43,7 +43,7 @@ static void mkl_fa_pack_q_fp16( dpct::queue_ptr stream, sycl::half * __restrict dst, const float * __restrict q_src, - int n_queries, int n_query_rows, int DKQ, + int n_queries, int DKQ, int gqa_ratio, int kvh_base_head, float q_scale, int64_t q_row_stride, int64_t q_head_stride, int64_t wg_size) { @@ -121,7 +121,7 @@ static void mkl_fa_online_softmax_chunk( float * __restrict VKQ_accum, int q0, int q_rows, int n_queries, int DV, int chunk_size, int chunk_start, - int kvh_head, int gqa_ratio, + int kvh_head, const sycl::half * mask_data, int64_t mask_head_stride, int64_t mask_row_stride, int mask_n_heads, float logit_softcap, int64_t wg_size) { @@ -473,7 +473,6 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * MKL_ACCUM(dequant_time_us, t_deq); // --- Resolve mask pointers --- - const sycl::half * mask_data = nullptr; int64_t mask_head_stride = 0; int64_t mask_row_stride = 0; int mask_n_heads = 0; @@ -547,7 +546,7 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * // 1. Pack all GQA Q heads into fp16 (full n_query_rows) mkl_fa_pack_q_fp16(stream, Q_head_f16_ptr, Q_batch, - n_queries, n_query_rows, DKQ, + n_queries, DKQ, gqa_ratio, kvh_base_head, q_scale, q_row_stride, q_head_stride, wg_size); @@ -605,7 +604,7 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * KQ_max_ptr, KQ_sum_ptr, VKQ_accum_ptr, q0, q_rows, n_queries, DV, this_chunk, chunk_start, - kvh_base_head, gqa_ratio, + kvh_base_head, mask_batch, mask_head_stride, mask_row_stride, mask_n_heads, logit_softcap, wg_size); diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 8372428795..7ebdce7fb9 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -921,16 +921,16 @@ ggml_backend_sycl_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft, void * dev_ptr; if (use_usm_system) { - GGML_SYCL_DEBUG("[SYCL] allocating %lu Bytes with USM system\n", size); + GGML_SYCL_DEBUG("[SYCL] allocating %zu Bytes with USM system\n", size); dev_ptr = (void *)aligned_malloc_host(alignment, aligned_size); if (!dev_ptr) { - GGML_LOG_ERROR("%s: can't allocate %lu Bytes of memory on host\n", __func__, size); + GGML_LOG_ERROR("%s: can't allocate %zu Bytes of memory on host\n", __func__, size); return nullptr; } } else { SYCL_CHECK(CHECK_TRY_ERROR(dev_ptr = (void *)ggml_sycl_malloc_device(size, *stream))); if (!dev_ptr) { - GGML_LOG_ERROR("%s: can't allocate %lu Bytes of memory on device\n", __func__, size); + GGML_LOG_ERROR("%s: can't allocate %zu Bytes of memory on device\n", __func__, size); return nullptr; } } @@ -1177,7 +1177,7 @@ ggml_backend_sycl_split_buffer_init_tensor(ggml_backend_buffer_t buffer, SYCL_CHECK(CHECK_TRY_ERROR(buf = (char *)ggml_sycl_malloc_device(size, *stream))); if (!buf) { char err_buf[1024]; - snprintf(err_buf, 1023, "%s: can't allocate %lu Bytes of memory on device\n", __func__, size); + snprintf(err_buf, 1023, "%s: can't allocate %zu Bytes of memory on device\n", __func__, size); throw std::runtime_error(err_buf); } // set padding to 0 to avoid possible NaN values @@ -1651,7 +1651,7 @@ struct ggml_sycl_pool_leg : public ggml_sycl_pool { SYCL_CHECK(CHECK_TRY_ERROR(ptr = (void *)ggml_sycl_malloc_device(look_ahead_size, *qptr))); if (!ptr) { - GGML_LOG_ERROR("%s: can't allocate %lu Bytes of memory on device/GPU\n", __func__, look_ahead_size); + GGML_LOG_ERROR("%s: can't allocate %zu Bytes of memory on device/GPU\n", __func__, look_ahead_size); return nullptr; } @@ -1663,7 +1663,7 @@ struct ggml_sycl_pool_leg : public ggml_sycl_pool { (uint32_t)(max_size/1024/1024), (uint32_t)(g_sycl_pool_size[id]/1024/1024), (uint32_t)(size/1024/1024)); #endif - // GGML_SYCL_DEBUG("ggml_sycl_pool_malloc_leg look_ahead_size=%lu, return %p\n", look_ahead_size, ptr); + // GGML_SYCL_DEBUG("ggml_sycl_pool_malloc_leg look_ahead_size=%zu, return %p\n", look_ahead_size, ptr); return ptr; } @@ -1843,7 +1843,7 @@ struct ggml_sycl_pool_host : public ggml_sycl_pool { SYCL_CHECK(CHECK_TRY_ERROR(ptr = (void *) sycl::malloc_host(size, *qptr))); if (!ptr) { - GGML_LOG_ERROR("%s: can't allocate %lu Bytes of memory on host\n", __func__, size); + GGML_LOG_ERROR("%s: can't allocate %zu Bytes of memory on host\n", __func__, size); return nullptr; } pool_size += size; @@ -2779,9 +2779,9 @@ inline void ggml_sycl_op_mul_mat_sycl( const float * src1_ddf1_i = src1->type == GGML_TYPE_F32 ? (const float *) src1_ddf_i : src1_ddq_as_f32.get(); { +#if GGML_SYCL_DNNL const int64_t gemm_flops = (int64_t)row_diff * src1_ncols * ne10; const bool use_mkl_direct = gemm_flops < 256 * 256 * 256; -#if GGML_SYCL_DNNL if (g_ggml_sycl_enable_dnn && !use_mkl_direct) { DnnlGemmWrapper::row_gemm(ctx, row_diff, src1_ncols, ne10, src0_ddf_i, DnnlGemmWrapper::to_dt(), src1_ddf1_i, DnnlGemmWrapper::to_dt(), @@ -3518,7 +3518,9 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, cons float * dst_ddf = static_cast(dst->data); const sycl::half * src1_f16 = static_cast(src1->data); +#if GGML_SYCL_DNNL const size_t type_size_src0 = ggml_type_size(src0->type); +#endif const size_t type_size_src1 = ggml_type_size(src1->type); bool is_src0_cont_2 = ggml_is_contiguous_2(src0); @@ -3535,6 +3537,7 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, cons scope_op_debug_print scope_dbg_print(__func__, "/to_fp16_nc_sycl", dst, /*num_src=*/2, " : converting src1 to fp16"); +#if GGML_SYCL_DNNL // iterate tensor dims and find the slowest moving dim and stride int last_dim=0; int last_str=0; @@ -3554,7 +3557,6 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, cons } } -#if GGML_SYCL_DNNL // oneDNN handles strided data and does not need overhead of ggml_get_to_fp16_nc_sycl const int64_t ne_src1 = src1->nb[last_str] * src1->ne[last_dim] / type_size_src1; src1_f16_alloc.alloc(ne_src1); diff --git a/ggml/src/ggml-sycl/im2col.cpp b/ggml/src/ggml-sycl/im2col.cpp index 7bf3584fb9..e666167594 100644 --- a/ggml/src/ggml-sycl/im2col.cpp +++ b/ggml/src/ggml-sycl/im2col.cpp @@ -85,7 +85,7 @@ static void im2col_sycl(const float * x, */ stream->parallel_for(sycl::nd_range<3>(block_nums * sycl::range<3>(1, 1, MIN(IC_KH_KW, SYCL_IM2COL_BLOCK_SIZE)), sycl::range<3>(1, 1, MIN(IC_KH_KW, SYCL_IM2COL_BLOCK_SIZE))), - [=](sycl::nd_item<3> item_ct1) { + [=](sycl::nd_item<3>) { im2col_kernel(x, dst, IC, IW, IH, OH, OW, KW, KH, IC_IH_IW, IH_IW, N_OH, KH_KW, IC_KH_KW, s0, s1, p0, p1, d0, d1); }); @@ -271,7 +271,7 @@ static void im2col_3d_sycl(const float * src, */ stream->parallel_for(sycl::nd_range<3>(block_nums * sycl::range<3>(1, 1, MIN(IC_KD_KH_KW, SYCL_IM2COL_BLOCK_SIZE)), sycl::range<3>(1, 1, MIN(IC_KD_KH_KW, SYCL_IM2COL_BLOCK_SIZE))), - [=](sycl::nd_item<3> item_ct1) { + [=](sycl::nd_item<3>) { im2col_3d_kernel(src, dst, N, IC, ID, IH, IW, OC, KD, KH, KW, OD, OH, OW, OH_OW, KD_KH_KW, ID_IH_IW, KH_KW, IH_IW, IC_ID_IH_IW, IC_KD_KH_KW, OW_KD_KH_KW, OD_OH_OW_IC_KD_KH_KW, OH_OW_IC_KD_KH_KW, OW_IC_KD_KH_KW, N_OD_OH, OD_OH, diff --git a/ggml/src/ggml-sycl/norm.cpp b/ggml/src/ggml-sycl/norm.cpp index 682a9f51ee..f98a7a9542 100644 --- a/ggml/src/ggml-sycl/norm.cpp +++ b/ggml/src/ggml-sycl/norm.cpp @@ -7,9 +7,6 @@ static void norm_f32(const float* x, float* dst, const int ncols, const int64_t dst_stride_col, const int64_t dst_stride_row, const int64_t dst_stride_channel, const int64_t dst_stride_sample, const float eps, const sycl::nd_item<3>& item_ct1, sycl::float2* s_sum, int block_size) { - const int nrows = item_ct1.get_group_range(2); - const int nchannels = item_ct1.get_group_range(1); - const int nthreads = item_ct1.get_local_range(2); const int sample = item_ct1.get_group(0); const int channel = item_ct1.get_group(1); @@ -155,9 +152,6 @@ static void rms_norm_f32(const float* x, float* dst, const int ncols, const float* mul = nullptr, const int64_t mul_stride_row = 0, const int64_t mul_stride_channel = 0, const int64_t mul_stride_sample = 0, const int mul_nrows = 0, const int mul_nchannels = 0, const int mul_nsamples = 0) { - const int nrows = item_ct1.get_group_range(2); - const int nchannels = item_ct1.get_group_range(1); - const int sample = item_ct1.get_group(0); const int channel = item_ct1.get_group(1); const int row = item_ct1.get_group(2); @@ -225,8 +219,6 @@ static void l2_norm_f32(const float * x, float * dst, const int ncols, const int64_t src_stride_sample, const int64_t dst_stride_col, const int64_t dst_stride_row, const int64_t dst_stride_channel, const int64_t dst_stride_sample, const float eps, const sycl::nd_item<3>& item_ct1, float* s_sum, const int block_size) { - const int nrows = item_ct1.get_group_range(2); - const int nchannels = item_ct1.get_group_range(1); const int row = item_ct1.get_group(2); const int channel = item_ct1.get_group(1); diff --git a/ggml/src/ggml-sycl/set_rows.cpp b/ggml/src/ggml-sycl/set_rows.cpp index 52a0bcb6eb..5f8d881a29 100644 --- a/ggml/src/ggml-sycl/set_rows.cpp +++ b/ggml/src/ggml-sycl/set_rows.cpp @@ -291,7 +291,7 @@ static void set_rows_sycl( stream->parallel_for( sycl::nd_range<1>(grid_size * block_size, block_size), - [=](sycl::nd_item<1> item_ct1) [[intel::reqd_sub_group_size(WARP_SIZE)]] { + [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { k_set_rows( src0_d, src1_d, dst_d, ne00, ne01, ne02, From 1cb3f5eb41d51fc98ac6b1d16ff199c427edd68d Mon Sep 17 00:00:00 2001 From: HumerousGorgon <31957201+HumerousGorgon@users.noreply.github.com> Date: Fri, 21 Aug 2026 13:16:29 +0800 Subject: [PATCH 07/38] sycl: Update gate logic for Alchemist GPUs regarding OneDNN features. (#26635) * feat: updated gating logic of fattn-onednn.cpp * verified device types * Update ggml/src/ggml-sycl/fattn-onednn.cpp Accepted recommendations to add bmg_g31 arch. Co-authored-by: Neo Zhang * Improved SPDA gate, added documentation. * Added arch var to reworked gate, fixing build errors. * Fix trailing whitespaces. --------- Co-authored-by: Neo Zhang --- ggml/src/ggml-sycl/fattn-onednn.cpp | 19 +++++++++++-------- 1 file changed, 11 insertions(+), 8 deletions(-) diff --git a/ggml/src/ggml-sycl/fattn-onednn.cpp b/ggml/src/ggml-sycl/fattn-onednn.cpp index fd17a25d5e..a501295192 100644 --- a/ggml/src/ggml-sycl/fattn-onednn.cpp +++ b/ggml/src/ggml-sycl/fattn-onednn.cpp @@ -21,14 +21,6 @@ bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst) { if (!g_ggml_sycl_fa_onednn) { return false; } - // Battlemage (Xe2) only, for now. On other Intel archs oneDNN's fused SDPA returns wrong results - // for some shapes (e.g. head_dim=64 on Arc / xe_hpg) -- an oneDNN bug tracked upstream at - // https://github.com/uxlfoundation/oneDNN/issues/5510. Remove this hardware limitation once that - // is fixed; until then non-BMG archs fall back to the existing FA kernel. - const gpu_arch arch = ggml_sycl_info().devices[ggml_sycl_get_device()].hw_info.arch; - if (arch != gpu_arch::intel_gpu_bmg_g21 && arch != gpu_arch::intel_gpu_bmg_g31) { - return false; - } const ggml_tensor * Q = dst->src[0]; const ggml_tensor * K = dst->src[1]; const ggml_tensor * V = dst->src[2]; @@ -60,6 +52,17 @@ bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst) { } } } + // This is the improved SPDA gate. Rather than gating Alchemist GPUs from all SPDA features, we instead target only the failing shapes. + // If the GPU being assessed isn't in the grouping below, it has full access to all SPDA shapes. Otherwise, if it's an Alchemist GPU, we block only the shapes with head sizes that fail. + // It is much easier to compare the device to a small list of failing cases than to define all the passing ones. + const gpu_arch arch = ggml_sycl_info().devices[ggml_sycl_get_device()].hw_info.arch; + bool support_spda = !(arch == gpu_arch::intel_gpu_dg2_g10 || + arch == gpu_arch::intel_gpu_dg2_g11 || + arch == gpu_arch::intel_gpu_dg2_g12); + + if (!support_spda && K->ne[0] == 64) { + return false; + } // Optional KV-length ceiling (GGML_SYCL_FA_ONEDNN_MAX_KV, 0 = unlimited). Escape hatch: // very long sequences make the fused SDPA slow enough to risk the xe driver watchdog on // some stacks; past the cap we fall back to the native FA kernel instead. From cd26896c19e6775b29a86908b5f049bbaec73305 Mon Sep 17 00:00:00 2001 From: Hongqiang Wang Date: Thu, 20 Aug 2026 22:30:17 -0700 Subject: [PATCH 08/38] opencl: keep the vocab-scale K-quant lm_head on the CPU for Adreno A7X (compiler issue workaround) (#26440) * opencl: keep the vocab-scale K-quant lm_head on the CPU on the Adreno A7X * opencl: revise comments --------- Co-authored-by: Li He --- ggml/src/ggml-opencl/ggml-opencl.cpp | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 26f952a173..84f854cc26 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -7393,6 +7393,19 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te op->src[0]->type == GGML_TYPE_Q4_K || op->src[0]->type == GGML_TYPE_Q5_K || op->src[0]->type == GGML_TYPE_Q6_K) { + // The E031.41 compiler (usually with A7x) miscompiles the flat K-quant + // GEMV kernels (kernel_mul_mv_q*_K_f32_flat) and makes lm_head run much + // slower than it should. So, make it fallback to CPU to preserve performance + // for this compiler series. + static const char * a7x_lmhead_env = getenv("GGML_OPENCL_A7X_LMHEAD_CPU"); + static const bool a7x_lmhead_cpu = (a7x_lmhead_env == nullptr || a7x_lmhead_env[0] != '0'); + if (a7x_lmhead_cpu && + backend_ctx->adreno_gen == ADRENO_GPU_GEN::A7X && + (op->src[0]->type == GGML_TYPE_Q4_K || op->src[0]->type == GGML_TYPE_Q5_K || + op->src[0]->type == GGML_TYPE_Q6_K) && + op->src[0]->ne[1] >= 32768) { // vocab-scale weight; no FFN/attn weight is this tall + return false; + } return op->src[1]->type == GGML_TYPE_F32 && ggml_is_contiguous(op->src[0]) && ggml_is_contiguous(op->src[1]); } else if (op->src[0]->type == GGML_TYPE_Q8_0) { return op->src[1]->type == GGML_TYPE_F32; From 9e89a196b8141c8bec0ad4c0bfd0cfc73bb8cdea Mon Sep 17 00:00:00 2001 From: Todd Malsbary Date: Fri, 21 Aug 2026 00:23:02 -0700 Subject: [PATCH 09/38] sycl : Add Q5_K ESIMD kernel (#26376) * Add DMMV Q4_K and Q6_K ESIMD kernels Configure cmake build with -DGGML_SYCL_ESIMD=ON to enable. Signed-off-by: Todd Malsbary * Refactor ESIMD kernels to share common code Signed-off-by: Todd Malsbary * Move control of ESIMD from compile to runtime Signed-off-by: Todd Malsbary * Use ESIMD by default when available Signed-off-by: Todd Malsbary * Fix possible error when using ESIMD by default While not an issue in the current version, this will become an issue when additional QK ESIMD kernels are added (such as Q2_K). Signed-off-by: Todd Malsbary * Add explicit unroll to ESIMD kernels Signed-off-by: Todd Malsbary * Tidy up ESIMD kernels a bit Signed-off-by: Todd Malsbary * Add DMMV Q5_K ESIMD kernel Signed-off-by: Todd Malsbary * Remove redundant copyright notice Signed-off-by: Todd Malsbary --------- Signed-off-by: Todd Malsbary --- ggml/src/ggml-sycl/dmmv.cpp | 27 ++++++- ggml/src/ggml-sycl/esimd.hpp | 134 ++++++++++++++++++++++++++++--- ggml/src/ggml-sycl/ggml-sycl.cpp | 1 + 3 files changed, 149 insertions(+), 13 deletions(-) diff --git a/ggml/src/ggml-sycl/dmmv.cpp b/ggml/src/ggml-sycl/dmmv.cpp index d8da0a16ba..fdcadbf91f 100644 --- a/ggml/src/ggml-sycl/dmmv.cpp +++ b/ggml/src/ggml-sycl/dmmv.cpp @@ -1955,6 +1955,23 @@ static void dequantize_mul_mat_vec_q4_K_sycl_reorder_esimd(const void *vx, const }); } +static void dequantize_mul_mat_vec_q5_K_sycl_reorder_esimd(const void *vx, const float *y, + float *dst, const int ncols, + const int nrows, + dpct::queue_ptr stream) { + GGML_ASSERT(ncols % QK_K == 0); + const int workgroups = (nrows + 1) / 2; + stream->submit([&](sycl::handler &h) { + sycl::local_accessor lmem(sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE * 2), h); + h.parallel_for( + sycl::nd_range<1>(sycl::range<1>((size_t)workgroups * GGML_SYCL_DMMV_ESIMD_WG_SIZE), sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE)), + [=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] { + dequantize_mul_mat_vec_reorder_esimd( + vx, y, dst, ncols, nrows, lmem, it); + }); + }); +} + static void dequantize_mul_mat_vec_q6_K_sycl_reorder_esimd(const void *vx, const float *y, float *dst, const int ncols, const int nrows, @@ -2134,7 +2151,15 @@ void ggml_sycl_op_dequantize_mul_mat_vec( case GGML_TYPE_Q5_K: if ((ggml_tensor_extra_gpu *) dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { - dequantize_mul_mat_vec_q5_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); +#ifdef GGML_SYCL_DMMV_HAS_ESIMD + if (g_ggml_sycl_enable_esimd) { + dequantize_mul_mat_vec_q5_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); + } + else +#endif + { + dequantize_mul_mat_vec_q5_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); + } } else { dequantize_mul_mat_vec_q5_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); } diff --git a/ggml/src/ggml-sycl/esimd.hpp b/ggml/src/ggml-sycl/esimd.hpp index d7609b11fe..04e596ed3d 100644 --- a/ggml/src/ggml-sycl/esimd.hpp +++ b/ggml/src/ggml-sycl/esimd.hpp @@ -1,15 +1,3 @@ -// -// MIT license -// Copyright (C) 2026 Intel Corporation -// SPDX-License-Identifier: MIT -// - -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// - #ifndef GGML_SYCL_ESIMD_HPP #define GGML_SYCL_ESIMD_HPP @@ -287,6 +275,128 @@ template <> struct esimd_reorder_q_traits { } }; +// --------------------------------------------------------------------------- +// Q5_K, SOA reorder layout produced by reorder_qw_q5_k: +// [qs: nb*(QK_K/2)] [qh: nb*(QK_K/8)] [scales: nb*K_SCALE_SIZE] [dm: nb*sizeof(half2)] +// with nb = nrows*num_blocks_per_row. +// +// Identical to Q4_K except each 4-bit quant gains a 5th (high) bit from qh: +// output chunk c (0..7) adds 16 when bit c of qh[l] is set, where qh[l] indexes +// the same 32 bytes for every chunk (matches dequantize_row_q5_K). +// --------------------------------------------------------------------------- +template <> struct esimd_reorder_q_traits { + struct ptrs { + const uint8_t * qs; + const uint8_t * qh; + const uint8_t * scales; + const sycl::half * dm; + }; + + static ESIMD_INLINE ptrs make_ptrs(const void * vx, size_t nb) { + const uint8_t * qs = (const uint8_t *) vx; + const uint8_t * qh = qs + nb * (QK_K / 2); + const uint8_t * scales = qh + nb * (QK_K / 8); + const sycl::half * dm = (const sycl::half *) (scales + nb * K_SCALE_SIZE); + return { qs, qh, scales, dm }; + } + + // extract bit `bit` (0..7) of each lane and move it to bit position 4, + // e.g. for the 4-bit base quant's 5th (high) bit. `bit` is always a + // compile-time-known unrolled loop constant at call sites, so this folds + // to a single mask (bit==4), mask+left-shift (bit<4), or mask+right-shift + // (bit>4) instead of the shift+mask+shift a naive `(qh>>bit & 1) << 4` emits. + static ESIMD_INLINE sycl::ext::intel::esimd::simd extract_bit_to_pos4( + sycl::ext::intel::esimd::simd qh, int bit) { + using namespace sycl::ext::intel::esimd; + simd masked = convert(qh & simd((uint8_t) (1u << bit))); + if (bit < 4) { + return masked << simd((uint16_t) (4 - bit)); + } else if (bit > 4) { + return masked >> simd((uint16_t) (bit - 4)); + } + return masked; + } + + static ESIMD_INLINE void mac_pair( + const ptrs & pa, size_t bia, + const ptrs & pb, size_t bib, bool has_b, + sycl::ext::intel::esimd::simd & y_vec, + sycl::ext::intel::esimd::simd & acc_a, + sycl::ext::intel::esimd::simd & acc_b) { + using namespace sycl::ext::intel::esimd; + + simd qs_a = block_load(pa.qs + bia * (QK_K / 2)); + simd qs_b = 0; + simd qh_a = block_load(pa.qh + bia * (QK_K / 8)); + simd qh_b = 0; + simd scales_a = block_load(pa.scales + bia * K_SCALE_SIZE); + simd scales_b = 0; + + const float dall_a = (float) pa.dm[bia * 2 + 0]; + const float dmin_a = (float) pa.dm[bia * 2 + 1]; + float dall_b = 0.0f; + float dmin_b = 0.0f; + if (has_b) { + qs_b = block_load(pb.qs + bib * (QK_K / 2)); + qh_b = block_load(pb.qh + bib * (QK_K / 8)); + scales_b = block_load(pb.scales + bib * K_SCALE_SIZE); + dall_b = (float) pb.dm[bib * 2 + 0]; + dmin_b = (float) pb.dm[bib * 2 + 1]; + } + + simd scale_f_a, min_f_a, scale_f_b, min_f_b; + unpack_scale_min_k4(scales_a, dall_a, dmin_a, scale_f_a, min_f_a); + unpack_scale_min_k4(scales_b, dall_b, dmin_b, scale_f_b, min_f_b); + + simd qs_lo_a = qs_a & simd(0x0F); + simd qs_hi_a = qs_a >> simd(4); + simd qs_lo_b = qs_b & simd(0x0F); + simd qs_hi_b = qs_b >> simd(4); + +#pragma unroll + for (int sb = 0; sb < 8; sb += 2) { + const int q_offset = sb * 16; + simd y_lo = y_vec.select<32, 1>(sb * 32); + simd y_hi = y_vec.select<32, 1>((sb + 1) * 32); + + const float scale_a_lo = scale_f_a[sb]; + const float scale_a_hi = scale_f_a[sb + 1]; + const float min_a_lo = min_f_a[sb]; + const float min_a_hi = min_f_a[sb + 1]; + const float scale_b_lo = scale_f_b[sb]; + const float scale_b_hi = scale_f_b[sb + 1]; + const float min_b_lo = min_f_b[sb]; + const float min_b_hi = min_f_b[sb + 1]; + + simd qa_lo_u8 = qs_lo_a.select<32, 1>(q_offset); + simd qa_hi_u8 = qs_hi_a.select<32, 1>(q_offset); + simd qb_lo_u8 = qs_lo_b.select<32, 1>(q_offset); + simd qb_hi_u8 = qs_hi_b.select<32, 1>(q_offset); + simd qa_lo = convert(qa_lo_u8); + simd qa_hi = convert(qa_hi_u8); + simd qb_lo = convert(qb_lo_u8); + simd qb_hi = convert(qb_hi_u8); + + // add the 5th bit: chunk sb uses qh bit sb, chunk sb+1 uses qh bit sb+1; + // qh always indexes the same 32 bytes regardless of chunk + qa_lo += extract_bit_to_pos4(qh_a, sb); + qa_hi += extract_bit_to_pos4(qh_a, sb + 1); + qb_lo += extract_bit_to_pos4(qh_b, sb); + qb_hi += extract_bit_to_pos4(qh_b, sb + 1); + + simd deq_a_lo = convert(qa_lo) * scale_a_lo + min_a_lo; + simd deq_a_hi = convert(qa_hi) * scale_a_hi + min_a_hi; + simd deq_b_lo = convert(qb_lo) * scale_b_lo + min_b_lo; + simd deq_b_hi = convert(qb_hi) * scale_b_hi + min_b_hi; + + acc_a += y_lo * deq_a_lo; + acc_b += y_lo * deq_b_lo; + acc_a += y_hi * deq_a_hi; + acc_b += y_hi * deq_b_hi; + } + } +}; + // --------------------------------------------------------------------------- // Q6_K, SOA reorder layout: // [ql: nb*(QK_K/2)] [qh: nb*(QK_K/4)] [scales(int8): nb*(QK_K/16)] [d: nb*half] diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 7ebdce7fb9..de56ea5b91 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -3811,6 +3811,7 @@ static bool ggml_sycl_supports_reorder_esimd(enum ggml_type type) { switch (type) { case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: + case GGML_TYPE_Q5_K: case GGML_TYPE_Q6_K: return true; default: From 5fff128451d7603857597ee1fc18ac1dfb90f148 Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Fri, 21 Aug 2026 10:29:17 +0300 Subject: [PATCH 10/38] test : make the FA V-is-view-of-K case a test case parameter (#27394) Resolve the TODO in test_flash_attn_ext: the branch that creates V as a sub-view of K (MLA-based models) was hardcoded for the 576/512 head shapes. Add a v_is_view_of_k test case parameter (default false) and select the sub-view branch on it; the existing 576/512 (DeepSeek MLA) cases now pass it explicitly, so the test coverage is unchanged. Also add more V-is-sub-view-of-K cases: the 320/256 (Mistral4 MLA) and 192/128 head shapes, and full views with equal head sizes (128/128 F16, 64/64 q8_0). Assisted-by: pi:llama.cpp/Qwen3.8-27B --- tests/test-backend-ops.cpp | 36 ++++++++++++++++++++++-------------- 1 file changed, 22 insertions(+), 14 deletions(-) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 89b954a7d1..8e3b273a1e 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -7085,9 +7085,10 @@ struct test_flash_attn_ext : public test_case { const ggml_type type_V; std::array permute; const bool kv_view; // create K/V as views of a larger buffer (like a KV cache) + const bool v_is_view_of_k; std::string vars() override { - return VARS_TO_STR15(hsk, hsv, nh, nr23, kv, nb, mask, sinks, max_bias, logit_softcap, prec, type_K, type_V, permute, kv_view); + return VARS_TO_STR16(hsk, hsv, nh, nr23, kv, nb, mask, sinks, max_bias, logit_softcap, prec, type_K, type_V, permute, kv_view, v_is_view_of_k); } double max_nmse_err() override { @@ -7104,9 +7105,9 @@ struct test_flash_attn_ext : public test_case { test_flash_attn_ext(int64_t hsk = 128, int64_t hsv = 128, int64_t nh = 32, std::array nr23 = {1, 1}, int64_t kv = 96, int64_t nb = 8, bool mask = true, bool sinks = false, float max_bias = 0.0f, float logit_softcap = 0.0f, ggml_prec prec = GGML_PREC_F32, ggml_type type_K = GGML_TYPE_F16, ggml_type type_V = GGML_TYPE_F16, std::array permute = {0, 1, 2, 3}, - bool kv_view = true) + bool kv_view = true, bool v_is_view_of_k = false) : hsk(hsk), hsv(hsv), nh(nh), nr23(nr23), kv(kv), nb(nb), mask(mask), sinks(sinks), max_bias(max_bias), logit_softcap(logit_softcap), prec(prec), - type_K(type_K), type_V(type_V), permute(permute), kv_view(kv_view) {} + type_K(type_K), type_V(type_V), permute(permute), kv_view(kv_view), v_is_view_of_k(v_is_view_of_k) {} ggml_tensor * build_graph(ggml_context * ctx) override { const int64_t hsk_padded = GGML_PAD(hsk, ggml_blck_size(type_K)); @@ -7138,14 +7139,14 @@ struct test_flash_attn_ext : public test_case { ggml_set_name(k, "k"); ggml_tensor * v = nullptr; - if (type_K == type_V && hsk_padded == 576 && hsv_padded == 512) { - // TODO: this branch should become a separate test case parameter instead of hardcoding this for these head shapes - - // in this branch, the V cache is sub-view of the K cache. this is used by some MLA-based models + if (v_is_view_of_k) { + // the V cache is a sub-view of the K cache. this is used by some MLA-based models // for more info: // - https://github.com/ggml-org/llama.cpp/pull/13435 // - https://github.com/ggml-org/llama.cpp/pull/18953#issuecomment-3774948392 // - https://github.com/ggml-org/llama.cpp/pull/18986 + GGML_ASSERT(type_K == type_V && hsv_padded <= hsk_padded); + v = ggml_view_4d(ctx, k, hsv_padded, kv, nh, nr23[1], k->nb[1], k->nb[2], k->nb[3], 0); } else { v = create_permuted(type_V, hsv_padded, kv, nh, nr23[1], kv_view); // the V tensor is usually a view of the V cache @@ -9906,12 +9907,14 @@ static std::vector> make_test_cases_eval() { if (hsk != 128 && prec == GGML_PREC_DEFAULT) continue; for (ggml_type type_KV : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_BF16, GGML_TYPE_Q8_0, GGML_TYPE_Q5_1, GGML_TYPE_Q5_0, GGML_TYPE_Q4_1, GGML_TYPE_Q4_0, GGML_TYPE_IQ4_NL}) { if (type_KV != GGML_TYPE_F16 && hsk != 64 && hsk != 72) continue; + // DeepSeek MLA: the V cache is a sub-view of the K cache + const bool v_is_view_of_k = hsk == 576; test_cases.emplace_back(new test_flash_attn_ext( - hsk, hsv, nh, {nr2, nr3}, kv, nb, mask, sinks, max_bias, logit_softcap, prec, type_KV, type_KV)); + hsk, hsv, nh, {nr2, nr3}, kv, nb, mask, sinks, max_bias, logit_softcap, prec, type_KV, type_KV, {0, 1, 2, 3}, true, v_is_view_of_k)); // run fewer test cases permuted if (mask == true && max_bias == 0.0f && logit_softcap == 0 && kv == 512) { test_cases.emplace_back(new test_flash_attn_ext( - hsk, hsv, nh, {nr2, nr3}, kv, nb, mask, sinks, max_bias, logit_softcap, prec, type_KV, type_KV, {0, 2, 1, 3})); + hsk, hsv, nh, {nr2, nr3}, kv, nb, mask, sinks, max_bias, logit_softcap, prec, type_KV, type_KV, {0, 2, 1, 3}, true, v_is_view_of_k)); } } } @@ -9950,11 +9953,16 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_flash_attn_ext(256, 256, 2, {16, 1}, 1025, 64, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 2, 1, 3})); test_cases.emplace_back(new test_flash_attn_ext(256, 256, 2, {16, 1}, 16384, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0)); - // MLA shape (V is a view of K) with quantized KV - // (the test harness builds V as a view of K for this shape; see build_graph) - test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {20, 1}, 113, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0)); - test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {20, 1}, 1024, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0)); - test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {20, 1}, 1024, 64, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0)); + // MLA shape: the V cache is a sub-view of the K cache, with quantized KV + test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {20, 1}, 113, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 1, 2, 3}, true, true)); + test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {20, 1}, 1024, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 1, 2, 3}, true, true)); + test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {20, 1}, 1024, 64, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 1, 2, 3}, true, true)); + + // more V-is-sub-view-of-K cases: other head shapes, and full views with equal head sizes + test_cases.emplace_back(new test_flash_attn_ext(320, 256, 1, {32, 1}, 512, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16, {0, 1, 2, 3}, true, true)); + test_cases.emplace_back(new test_flash_attn_ext(192, 128, 4, {8, 1}, 512, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16, {0, 1, 2, 3}, true, true)); + test_cases.emplace_back(new test_flash_attn_ext(128, 128, 8, {4, 1}, 512, 8, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16, {0, 1, 2, 3}, true, true)); + test_cases.emplace_back(new test_flash_attn_ext(64, 64, 4, {1, 1}, 512, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 1, 2, 3}, true, true)); // large-KV F16 cases (Qwen3.6-27B geometry and a llama-class control): the upstream matrix // stops at kv=1024, blind to long-context FA bugs (e.g. the oneDNN SDPA ordering race on BMG). From ff14356e0caf6988f61f1f15f9dfe7d5ab398271 Mon Sep 17 00:00:00 2001 From: Todd Malsbary Date: Fri, 21 Aug 2026 01:01:40 -0700 Subject: [PATCH 11/38] sycl : add Q2_K reordered MMVQ and ESIMD kernels (#26336) * Add DMMV Q4_K and Q6_K ESIMD kernels Configure cmake build with -DGGML_SYCL_ESIMD=ON to enable. Signed-off-by: Todd Malsbary * Refactor ESIMD kernels to share common code Signed-off-by: Todd Malsbary * Move control of ESIMD from compile to runtime Signed-off-by: Todd Malsbary * Use ESIMD by default when available Signed-off-by: Todd Malsbary * Fix possible error when using ESIMD by default While not an issue in the current version, this will become an issue when additional QK ESIMD kernels are added (such as Q2_K). Signed-off-by: Todd Malsbary * Add explicit unroll to ESIMD kernels Signed-off-by: Todd Malsbary * Tidy up ESIMD kernels a bit Signed-off-by: Todd Malsbary * Add a reordered Q2_K MMVQ kernel Signed-off-by: Todd Malsbary * Add DMMV Q2_K ESIMD kernel Signed-off-by: Todd Malsbary --------- Signed-off-by: Todd Malsbary --- ggml/src/ggml-sycl/convert.cpp | 25 ++++++++- ggml/src/ggml-sycl/dequantize.hpp | 41 +++++++++++++++ ggml/src/ggml-sycl/dmmv.cpp | 27 +++++++++- ggml/src/ggml-sycl/esimd.hpp | 87 +++++++++++++++++++++++++++++++ ggml/src/ggml-sycl/ggml-sycl.cpp | 2 + ggml/src/ggml-sycl/mmvq.cpp | 75 +++++++++++++++++++++++++- ggml/src/ggml-sycl/quants.hpp | 23 ++++++++ ggml/src/ggml-sycl/vecdotq.hpp | 33 ++++++++++++ 8 files changed, 309 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-sycl/convert.cpp b/ggml/src/ggml-sycl/convert.cpp index 9ec9276952..b660b56ab3 100644 --- a/ggml/src/ggml-sycl/convert.cpp +++ b/ggml/src/ggml-sycl/convert.cpp @@ -76,6 +76,19 @@ static void dequantize_row_q2_K_sycl(const void *vx, dst_t *y, const int64_t k, #endif } +template +static void dequantize_row_q2_K_sycl_reorder(const void *vx, dst_t *y, const int64_t k, + dpct::queue_ptr stream) { + const int64_t nb = k / QK_K; + + dpct::has_capability_or_fail(stream->get_device(), { sycl::aspect::fp16 }); + stream->parallel_for( + sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)), + [=](sycl::nd_item<3> item_ct1) { + dequantize_block_q2_K_reorder(vx, y, item_ct1, nb); + }); +} + template static void dequantize_row_q3_K_sycl(const void *vx, dst_t *y, const int64_t k, dpct::queue_ptr stream) { @@ -667,7 +680,11 @@ to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst) { return dequantize_block_sycl; } case GGML_TYPE_Q2_K: - return dequantize_row_q2_K_sycl; + if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { + return dequantize_row_q2_K_sycl_reorder; + } else { + return dequantize_row_q2_K_sycl; + } case GGML_TYPE_Q3_K: if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { return dequantize_row_q3_K_sycl_reorder; @@ -753,7 +770,11 @@ to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor *dst) { return dequantize_block_sycl; } case GGML_TYPE_Q2_K: - return dequantize_row_q2_K_sycl; + if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { + return dequantize_row_q2_K_sycl_reorder; + } else { + return dequantize_row_q2_K_sycl; + } case GGML_TYPE_Q3_K: if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { return dequantize_row_q3_K_sycl_reorder; diff --git a/ggml/src/ggml-sycl/dequantize.hpp b/ggml/src/ggml-sycl/dequantize.hpp index 876ba1b444..1b13e0f1a3 100644 --- a/ggml/src/ggml-sycl/dequantize.hpp +++ b/ggml/src/ggml-sycl/dequantize.hpp @@ -943,6 +943,47 @@ static void dequantize_block_q2_K(const void * __restrict__ vx, dst_t * __restri } +template +static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, + const sycl::nd_item<3> & item_ct1, int64_t n_blocks) { +#if QK_K == 256 + const int64_t i = item_ct1.get_group(2); + if (i >= n_blocks) { + return; + } + + const uint8_t * base = static_cast(vx); + const size_t qs_offset = i * (QK_K / 4); + const size_t scales_offset = n_blocks * (QK_K / 4) + i * (QK_K / 16); + const size_t dm_offset = n_blocks * (QK_K / 4) + n_blocks * (QK_K / 16) + i * sizeof(ggml_half2); + + const uint8_t * qs = base + qs_offset; + const uint8_t * scales = base + scales_offset; + const ggml_half2 * dm = reinterpret_cast(base + dm_offset); + + const int64_t tid = item_ct1.get_local_id(2); + const int64_t n = tid / 32; + const int64_t l = tid - 32 * n; + const int64_t is = 8 * n + l / 16; + + const uint8_t q = qs[32 * n + l]; + dst_t * y = yy + i * QK_K + 128 * n; + + const float dall = (*dm)[0]; + const float dmin = (*dm)[1]; + y[l+ 0] = dall * (scales[is+0] & 0xF) * ((q >> 0) & 3) - dmin * (scales[is+0] >> 4); + y[l+32] = dall * (scales[is+2] & 0xF) * ((q >> 2) & 3) - dmin * (scales[is+2] >> 4); + y[l+64] = dall * (scales[is+4] & 0xF) * ((q >> 4) & 3) - dmin * (scales[is+4] >> 4); + y[l+96] = dall * (scales[is+6] & 0xF) * ((q >> 6) & 3) - dmin * (scales[is+6] >> 4); +#else + GGML_UNUSED(vx); + GGML_UNUSED(yy); + GGML_UNUSED(item_ct1); + GGML_UNUSED(n_blocks); + GGML_ABORT("Q2_K reorder dequantize not supported for QK_K != 256"); +#endif +} + template static void dequantize_block_q3_K(const void * __restrict__ vx, dst_t * __restrict__ yy, const sycl::nd_item<3> &item_ct1) { diff --git a/ggml/src/ggml-sycl/dmmv.cpp b/ggml/src/ggml-sycl/dmmv.cpp index fdcadbf91f..d47d6831a3 100644 --- a/ggml/src/ggml-sycl/dmmv.cpp +++ b/ggml/src/ggml-sycl/dmmv.cpp @@ -1921,6 +1921,23 @@ ESIMD_INLINE void dequantize_mul_mat_vec_reorder_esimd( } } +static void dequantize_mul_mat_vec_q2_K_sycl_reorder_esimd(const void *vx, const float *y, + float *dst, const int ncols, + const int nrows, + dpct::queue_ptr stream) { + GGML_ASSERT(ncols % QK_K == 0); + const int workgroups = (nrows + 1) / 2; + stream->submit([&](sycl::handler &h) { + sycl::local_accessor lmem(sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE * 2), h); + h.parallel_for( + sycl::nd_range<1>(sycl::range<1>((size_t)workgroups * GGML_SYCL_DMMV_ESIMD_WG_SIZE), sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE)), + [=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] { + dequantize_mul_mat_vec_reorder_esimd( + vx, y, dst, ncols, nrows, lmem, it); + }); + }); +} + static void dequantize_mul_mat_vec_q3_K_sycl_reorder_esimd(const void *vx, const float *y, float *dst, const int ncols, const int nrows, @@ -2111,7 +2128,15 @@ void ggml_sycl_op_dequantize_mul_mat_vec( case GGML_TYPE_Q2_K: if ((ggml_tensor_extra_gpu *) dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { - dequantize_mul_mat_vec_q2_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); +#ifdef GGML_SYCL_DMMV_HAS_ESIMD + if (g_ggml_sycl_enable_esimd) { + dequantize_mul_mat_vec_q2_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); + } + else +#endif + { + dequantize_mul_mat_vec_q2_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); + } } else { dequantize_mul_mat_vec_q2_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); } diff --git a/ggml/src/ggml-sycl/esimd.hpp b/ggml/src/ggml-sycl/esimd.hpp index 04e596ed3d..0485ff0cee 100644 --- a/ggml/src/ggml-sycl/esimd.hpp +++ b/ggml/src/ggml-sycl/esimd.hpp @@ -61,6 +61,93 @@ static ESIMD_INLINE void unpack_scale_min_k4( min_f = convert(m) * (-dmin); } +// --------------------------------------------------------------------------- +// Q2_K, SOA reorder layout produced by reorder_qw_q2_k: +// [qs: nb*(QK_K/4)] [scales: nb*(QK_K/16)] [dm: nb*sizeof(half2)] +// with nb = nrows*num_blocks_per_row. +// +// 2 bits per weight. The 8 output chunks of 32 (matching dequantize_row_q2_K) +// map to super-chunk s (0..7): byte base 32*(s/4) into the 64-byte qs array, +// bit shift 2*(s%4); the low 16 lanes use scales[2s], the high 16 use +// scales[2s+1], with dl = d*(sc & 0xF), ml = dmin*(sc >> 4), deq = dl*q - ml. +// --------------------------------------------------------------------------- +template <> struct esimd_reorder_q_traits { + struct ptrs { + const uint8_t * qs; + const uint8_t * scales; + const sycl::half * dm; + }; + + static ESIMD_INLINE ptrs make_ptrs(const void * vx, size_t nb) { + const uint8_t * qs = (const uint8_t *) vx; + const uint8_t * scales = qs + nb * (QK_K / 4); + const sycl::half * dm = (const sycl::half *) (scales + nb * (QK_K / 16)); + return { qs, scales, dm }; + } + + static ESIMD_INLINE void mac_pair( + const ptrs & pa, size_t bia, + const ptrs & pb, size_t bib, bool has_b, + sycl::ext::intel::esimd::simd & y_vec, + sycl::ext::intel::esimd::simd & acc_a, + sycl::ext::intel::esimd::simd & acc_b) { + using namespace sycl::ext::intel::esimd; + + simd qs_a = block_load(pa.qs + bia * (QK_K / 4)); + simd qs_b = 0; + simd scales_a = block_load(pa.scales + bia * (QK_K / 16)); + simd scales_b = 0; + + const float dall_a = (float) pa.dm[bia * 2 + 0]; + const float dmin_a = (float) pa.dm[bia * 2 + 1]; + float dall_b = 0.0f; + float dmin_b = 0.0f; + if (has_b) { + qs_b = block_load(pb.qs + bib * (QK_K / 4)); + scales_b = block_load(pb.scales + bib * (QK_K / 16)); + dall_b = (float) pb.dm[bib * 2 + 0]; + dmin_b = (float) pb.dm[bib * 2 + 1]; + } + + // per-chunk scale (d * (sc & 0xF)) and min (-dmin * (sc >> 4)), all 16 codes; + // min carries the negation so the dequant epilogue adds (matches Q4_K/Q5_K) + simd scale_f_a = convert(scales_a & simd(0x0F)) * dall_a; + simd min_f_a = convert(scales_a >> simd(4)) * (-dmin_a); + simd scale_f_b = convert(scales_b & simd(0x0F)) * dall_b; + simd min_f_b = convert(scales_b >> simd(4)) * (-dmin_b); + +#pragma unroll + for (int s = 0; s < 8; ++s) { + const int byte_base = 32 * (s / 4); + const uint8_t shift = (uint8_t) (2 * (s % 4)); + simd y_s = y_vec.select<32, 1>(s * 32); + + simd qa = (qs_a.select<32, 1>(byte_base) >> shift) & simd(3); + simd qb = (qs_b.select<32, 1>(byte_base) >> shift) & simd(3); + + const float scale_a_lo = scale_f_a[2 * s + 0]; + const float scale_a_hi = scale_f_a[2 * s + 1]; + const float min_a_lo = min_f_a[2 * s + 0]; + const float min_a_hi = min_f_a[2 * s + 1]; + const float scale_b_lo = scale_f_b[2 * s + 0]; + const float scale_b_hi = scale_f_b[2 * s + 1]; + const float min_b_lo = min_f_b[2 * s + 0]; + const float min_b_hi = min_f_b[2 * s + 1]; + + simd scale_vec_a = splat_lo_hi(scale_a_lo, scale_a_hi); + simd min_vec_a = splat_lo_hi(min_a_lo, min_a_hi); + simd scale_vec_b = splat_lo_hi(scale_b_lo, scale_b_hi); + simd min_vec_b = splat_lo_hi(min_b_lo, min_b_hi); + + simd deq_a = convert(qa) * scale_vec_a + min_vec_a; + simd deq_b = convert(qb) * scale_vec_b + min_vec_b; + + acc_a += y_s * deq_a; + acc_b += y_s * deq_b; + } + } +}; + // --------------------------------------------------------------------------- // Q3_K, SOA reorder layout produced by reorder_qw_q3_k: // [qs: nb*(QK_K/4)] [hmask: nb*(QK_K/8)] [scales: nb*12] [d: nb*sizeof(half)] diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index de56ea5b91..3f82020f41 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -3796,6 +3796,7 @@ inline bool ggml_sycl_supports_reorder_mmvq(enum ggml_type type) { case GGML_TYPE_Q1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q8_0: + case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: case GGML_TYPE_Q5_K: @@ -3809,6 +3810,7 @@ inline bool ggml_sycl_supports_reorder_mmvq(enum ggml_type type) { static bool ggml_sycl_supports_reorder_esimd(enum ggml_type type) { #ifdef GGML_SYCL_DMMV_HAS_ESIMD switch (type) { + case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: case GGML_TYPE_Q5_K: diff --git a/ggml/src/ggml-sycl/mmvq.cpp b/ggml/src/ggml-sycl/mmvq.cpp index 123b2a2f03..bfccb4b08f 100644 --- a/ggml/src/ggml-sycl/mmvq.cpp +++ b/ggml/src/ggml-sycl/mmvq.cpp @@ -1401,6 +1401,64 @@ static void mul_mat_vec_q2_K_q8_1_sycl_switch_ncols( } } +static void reorder_mul_mat_vec_q2_k_q8_1_sycl(const void * vx, const void * vy, float * dst, const int ncols, + const int nrows, dpct::queue_ptr stream) { + GGML_ASSERT(ncols % QK_K == 0); + + // Round up to a whole number of subgroup-sized workgroups; out-of-range rows are skipped inside the kernel. + constexpr size_t num_subgroups = WARP_SIZE; + const int block_num_y = ceil_div(nrows, GGML_SYCL_MMV_Y * (int) num_subgroups); + const sycl::range<3> block_nums(1, 1, block_num_y); + const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE); + + stream->submit([&](sycl::handler & cgh) { + cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), + [=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { + mul_mat_vec_q_reorder>(vx, vy, dst, ncols, nrows, + nd_item); + }); + }); +} + +template +static void reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols( + const void * vx, const void * vy, float * dst, + const int ncols, const int nrows, + const int stride_col_y_bytes, const int stride_col_dst, + dpct::queue_ptr stream) { + GGML_ASSERT(ncols % QK_K == 0); + constexpr size_t num_subgroups = WARP_SIZE; + const int block_num_y = ceil_div(nrows, GGML_SYCL_MMV_Y * (int) num_subgroups); + const sycl::range<3> block_nums(1, 1, block_num_y); + const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE); + + stream->submit([&](sycl::handler & cgh) { + cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), + [=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { + mul_mat_vec_q_reorder_ncols, ncols_dst>( + vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, nd_item); + }); + }); +} + +static void reorder_mul_mat_vec_q2_k_q8_1_sycl_switch_ncols( + const void * vx, const void * vy, float * dst, + const int ncols, const int nrows, const int ncols_dst, + const int stride_col_y_bytes, const int stride_col_dst, + dpct::queue_ptr stream) { + switch (ncols_dst) { + case 1: reorder_mul_mat_vec_q2_k_q8_1_sycl(vx, vy, dst, ncols, nrows, stream); break; + case 2: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<2>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; + case 3: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<3>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; + case 4: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<4>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; + case 5: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<5>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; + case 6: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<6>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; + case 7: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<7>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; + case 8: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<8>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; + default: GGML_ABORT("unsupported ncols_dst=%d for Q2_K reorder multi-col MMVQ", ncols_dst); + } +} + static void mul_mat_vec_q3_K_q8_1_sycl(const void *vx, const void *vy, float *dst, const int ncols, const int nrows, @@ -2297,7 +2355,21 @@ void ggml_sycl_op_mul_mat_vec_q(ggml_backend_sycl_context & ctx, const ggml_tens } break; case GGML_TYPE_Q2_K: - if (i == 0 && src1_ncols > 1 && src1_ncols <= 8) { + if ((ggml_tensor_extra_gpu *) dst->src[0]->extra && + ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { + if (i == 0 && src1_ncols > 1 && src1_ncols <= 8) { + const int stride_col_y_bytes = src1_padded_col_size * q8_1_ts / q8_1_bs; + const int stride_col_dst = dst->ne[0]; + GGML_SYCL_DEBUG("Calling reorder_mul_mat_vec_q2_k_q8_1_sycl_switch_ncols ncols=%d\n", (int)src1_ncols); + reorder_mul_mat_vec_q2_k_q8_1_sycl_switch_ncols( + src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, + src1_ncols, stride_col_y_bytes, stride_col_dst, stream); + return; + } else { + GGML_SYCL_DEBUG("Calling reorder_mul_mat_vec_q2_k_q8_1_sycl\n"); + reorder_mul_mat_vec_q2_k_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream); + } + } else if (i == 0 && src1_ncols > 1 && src1_ncols <= 8) { const int stride_col_y = src1_padded_col_size / QK8_1; const int stride_col_dst = dst->ne[0]; GGML_SYCL_DEBUG("Calling mul_mat_vec_q2_K_q8_1_sycl_switch_ncols ncols=%d\n", (int)src1_ncols); @@ -2306,6 +2378,7 @@ void ggml_sycl_op_mul_mat_vec_q(ggml_backend_sycl_context & ctx, const ggml_tens src1_ncols, stride_col_y, stride_col_dst, stream); return; } else if (i == 0 || src1_ncols == 1) { + GGML_SYCL_DEBUG("Calling mul_mat_vec_q2_K_q8_1_sycl\n"); mul_mat_vec_q2_K_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream); } break; diff --git a/ggml/src/ggml-sycl/quants.hpp b/ggml/src/ggml-sycl/quants.hpp index 95287f1751..a26a6ce6e6 100644 --- a/ggml/src/ggml-sycl/quants.hpp +++ b/ggml/src/ggml-sycl/quants.hpp @@ -58,6 +58,29 @@ template <> struct block_q_t { static constexpr int block_to_q8_1_ratio() { return traits::qk / QK8_1; } }; +template <> struct block_q_t { + struct traits { + static constexpr uint32_t qk = QK_K; + static constexpr uint32_t qi = QI2_K; + static constexpr uint32_t qr = QR2_K; + static constexpr uint32_t vdr_mmvq = 1; + }; + + // Reordered layout: [qs (QK_K/4 per block)] [scales (QK_K/16 per block)] [dm] + static constexpr std::pair get_block_offset(const int block_index, const int /* n_blocks */) { + return { block_index * (QK_K / 4), 0 }; + } + + static constexpr std::pair get_d_offset(int nrows, int ncols, const int block_index) { + auto nblocks = (nrows * (ncols / QK_K)); + auto total_qs_bytes = nblocks * (QK_K / 4); + return { total_qs_bytes + block_index * (QK_K / 16), + total_qs_bytes + nblocks * (QK_K / 16) + block_index * sizeof(ggml_half2) }; + } + + static constexpr int block_to_q8_1_ratio() { return traits::qk / QK8_1; } +}; + template <> struct block_q_t { struct traits { static constexpr uint32_t qk = QK_K; diff --git a/ggml/src/ggml-sycl/vecdotq.hpp b/ggml/src/ggml-sycl/vecdotq.hpp index c11a6e8f9c..3ad4cee93a 100644 --- a/ggml/src/ggml-sycl/vecdotq.hpp +++ b/ggml/src/ggml-sycl/vecdotq.hpp @@ -429,6 +429,39 @@ template <> struct reorder_vec_dot_q_sycl { } }; +template <> struct reorder_vec_dot_q_sycl { + static constexpr ggml_type gtype = GGML_TYPE_Q2_K; + + using q2_k_block = ggml_sycl_reordered::block_q_t; + using q2_k_traits = typename q2_k_block::traits; + + __dpct_inline__ float operator()(const void * __restrict__ vbq, const std::pair ibx_offset, + const std::pair d_offset, const int8_t * q8_1_quant_ptr, + const sycl::half2 * q8_1_ds, const int & iqs) { + const uint8_t * base = static_cast(vbq); + const uint8_t * qs = base + ibx_offset.first; + const uint8_t * scales = base + d_offset.first; + const ggml_half2 * dm = reinterpret_cast(base + d_offset.second); + + const int bq8_offset = QR2_K * (iqs / QI8_1); + const int scale_offset = iqs - iqs % QI8_1 + (iqs % QI8_1) / (QI8_1 / 2); + + const int v = get_int_from_uint8_aligned(qs, iqs); + + int u[QR2_K]; + float d8[QR2_K]; + +#pragma unroll + for (int i = 0; i < QR2_K; ++i) { + const int8_t * quant_base_ptr = q8_1_quant_ptr + (bq8_offset + i) * QK8_1; + u[i] = get_int_from_int8_aligned(quant_base_ptr, iqs % QI8_1); + d8[i] = (*(q8_1_ds + bq8_offset + i))[0]; + } + + return vec_dot_q2_K_q8_1_impl_mmvq(v, u, scales + scale_offset, *dm, d8); + } +}; + template <> struct reorder_vec_dot_q_sycl { static constexpr ggml_type gtype = GGML_TYPE_Q3_K; From 62b22690602665a0f34ce5f722915839f4e4913d Mon Sep 17 00:00:00 2001 From: Charles Xu Date: Fri, 21 Aug 2026 10:33:30 +0200 Subject: [PATCH 12/38] kleidiai : add SME2 F32 GEMV kernel support (#26891) --- ggml/src/ggml-cpu/CMakeLists.txt | 3 ++ ggml/src/ggml-cpu/kleidiai/kernels.cpp | 48 ++++++++++++++++--------- ggml/src/ggml-cpu/kleidiai/kleidiai.cpp | 48 ++++++++++++++++++------- 3 files changed, 70 insertions(+), 29 deletions(-) diff --git a/ggml/src/ggml-cpu/CMakeLists.txt b/ggml/src/ggml-cpu/CMakeLists.txt index 32e1e7aa19..e16ac996a4 100644 --- a/ggml/src/ggml-cpu/CMakeLists.txt +++ b/ggml/src/ggml-cpu/CMakeLists.txt @@ -639,6 +639,7 @@ function(ggml_add_cpu_backend_variant_impl tag_name) ${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_fp32_bf16p_bf16p/ ${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_f32_f16p_qsi4c32p/ ${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_f32_f32p_f32p/ + ${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_f32_f32_f32p/ ${KLEIDIAI_SRC}/kai/ukernels/matmul/pack/) set(ARCH_FLAGS_TEMP "${ARCH_FLAGS}") @@ -701,6 +702,8 @@ function(ggml_add_cpu_backend_variant_impl tag_name) ${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_f32_f16p_qsi4c32p/kai_matmul_clamp_f32_f16p1vlx2_qsi4c32p4vlx2_1vlx4vl_sme2_mopa_asm.S ${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_f32_f32p_f32p/kai_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa.c ${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_f32_f32p_f32p/kai_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa_asm.S + ${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_f32_f32_f32p/kai_matmul_clamp_f32_f32_f32p2vlx1b_1x16vl_sme2_mla.c + ${KLEIDIAI_SRC}/kai/ukernels/matmul/matmul_clamp_f32_f32_f32p/kai_matmul_clamp_f32_f32_f32p2vlx1b_1x16vl_sme2_mla_asm.S ${KLEIDIAI_SRC}/kai/ukernels/matmul/pack/kai_lhs_pack_bf16p2vlx2_f32_sme.c ${KLEIDIAI_SRC}/kai/ukernels/matmul/pack/kai_rhs_pack_kxn_bf16p2vlx2b_f32_x32_sme.c ${KLEIDIAI_SRC}/kai/ukernels/matmul/pack/kai_lhs_pack_f16pmrx2_f32_neon.c diff --git a/ggml/src/ggml-cpu/kleidiai/kernels.cpp b/ggml/src/ggml-cpu/kleidiai/kernels.cpp index 3c31ab9d35..70b519f29c 100644 --- a/ggml/src/ggml-cpu/kleidiai/kernels.cpp +++ b/ggml/src/ggml-cpu/kleidiai/kernels.cpp @@ -23,6 +23,7 @@ #include "kai_matmul_clamp_f32_qsi8d32p1x8_qsi4c32p8x8_1x8_sve_dotprod.h" #include "kai_matmul_clamp_f32_f16p1vlx2_qsi4c32p4vlx2_1vlx4vl_sme2_mopa.h" #include "kai_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa.h" +#include "kai_matmul_clamp_f32_f32_f32p2vlx1b_1x16vl_sme2_mla.h" #include "kai_matmul_clamp_f32_f32p2vlx1_f32p2vlx1b_2vlx2vl_sme_mopa.h" #include "kai_lhs_pack_bf16p2vlx2_f32_sme.h" @@ -76,6 +77,21 @@ static inline void kernel_run_fn10(size_t m, size_t n, size_t k, size_t /*bl*/, Fn(m, n, k, lhs, rhs, dst, dst_stride_row, dst_stride_col, clamp_min, clamp_max); } +template +static inline void kernel_run_lhs_stride_fn10(size_t m, + size_t n, + size_t k, + size_t lhs_stride, + const void * lhs, + const void * rhs, + void * dst, + size_t dst_stride_row, + size_t dst_stride_col, + float clamp_min, + float clamp_max) { + Fn(m, n, k, lhs, lhs_stride, rhs, dst, dst_stride_row, dst_stride_col, clamp_min, clamp_max); +} + template static inline void kernel_run_float_fn10(size_t m, size_t n, size_t k, size_t /*bl*/, const void* lhs, const void* rhs, void* dst, @@ -947,25 +963,25 @@ static ggml_kleidiai_kernels ggml_kleidiai_kernels_f32[] = { /* .packed_size_ex = */ &lhs_ps_fn5, /* .pack_func_ex = */ &lhs_pack_void_fn9, }, - /* SME GEMV */ + /* SME2 GEMV */ { - /* .get_m_step = */ kai_get_m_step_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa, - /* .get_n_step = */ kai_get_n_step_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa, - /* .get_mr = */ kai_get_mr_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa, - /* .get_nr = */ kai_get_nr_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa, - /* .get_kr = */ kai_get_kr_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa, - /* .get_sr = */ kai_get_sr_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa, - /* .get_dst_offset = */ kai_get_dst_offset_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa, - /* .get_dst_size = */ kai_get_dst_size_matmul_clamp_f32_f32p2vlx1_f32p2vlx1biasf32_sme2_mopa, - /* .get_lhs_offset_ex = */ nullptr, - /* .get_rhs_packed_offset_ex = */ nullptr, - /* .run_kernel_ex = */ nullptr, + /* .get_m_step = */ kai_get_m_step_matmul_clamp_f32_f32_f32p2vlx1b_1x16vl_sme2_mla, + /* .get_n_step = */ kai_get_n_step_matmul_clamp_f32_f32_f32p2vlx1b_1x16vl_sme2_mla, + /* .get_mr = */ kai_get_m_step_matmul_clamp_f32_f32_f32p2vlx1b_1x16vl_sme2_mla, + /* .get_nr = */ kai_get_nr_matmul_clamp_f32_f32_f32p2vlx1b_1x16vl_sme2_mla, + /* .get_kr = */ kai_get_kr_matmul_clamp_f32_f32_f32p2vlx1b_1x16vl_sme2_mla, + /* .get_sr = */ kai_get_sr_matmul_clamp_f32_f32_f32p2vlx1b_1x16vl_sme2_mla, + /* .get_dst_offset = */ kai_get_dst_offset_matmul_clamp_f32_f32_f32p2vlx1b_1x16vl_sme2_mla, + /* .get_dst_size = */ kai_get_dst_size_matmul_clamp_f32_f32_f32p2vlx1b_1x16vl_sme2_mla, + /* .get_lhs_offset_ex = */ &kernel_offs_fn2, + /* .get_rhs_packed_offset_ex = */ &kernel_offs_fn2, + /* .run_kernel_ex = */ &kernel_run_lhs_stride_fn10, }, /* .gemv_lhs_info = */ { - /* .get_offset = */ kai_get_lhs_offset_lhs_pack_f32p2vlx1_f32_sme, - /* .get_packed_offset_ex = */ &lhs_offs_fn5, - /* .packed_size_ex = */ &lhs_ps_fn5, - /* .pack_func_ex = */ &lhs_pack_void_fn9, + /* .get_offset = */ nullptr, + /* .get_packed_offset_ex = */ nullptr, + /* .packed_size_ex = */ nullptr, + /* .pack_func_ex = */ nullptr, }, /* .rhs_info = */ { /* .packed_stride = */ nullptr, diff --git a/ggml/src/ggml-cpu/kleidiai/kleidiai.cpp b/ggml/src/ggml-cpu/kleidiai/kleidiai.cpp index 2266c16898..6729ae8422 100644 --- a/ggml/src/ggml-cpu/kleidiai/kleidiai.cpp +++ b/ggml/src/ggml-cpu/kleidiai/kleidiai.cpp @@ -696,6 +696,15 @@ class tensor_traits : public ggml::cpu::tensor_traits { } if (op->src[0]->type == GGML_TYPE_F32) { + ggml_kleidiai_kernels * primary = kernel_chain[0]; + kernel_info * gemv_kernel = primary ? &primary->gemv : nullptr; + if (is_gemv && op->src[1]->nb[0] == (int64_t) sizeof(float) && gemv_kernel && + gemv_kernel->get_lhs_offset_ex && gemv_kernel->get_rhs_packed_offset_ex && + gemv_kernel->run_kernel_ex && gemv_kernel->get_dst_offset) { + size = 0; + return true; + } + size_t cursor = 0; bool any_slot = false; @@ -811,15 +820,28 @@ class tensor_traits : public ggml::cpu::tensor_traits { return false; } - kernel_info * kernel = &kernels->gemm; + const size_t k = ne00; + const size_t m = ne11; + const size_t n = ne01; + const bool use_gemv = m == 1 && src1->nb[0] == (int64_t) sizeof(float) && + kernels->gemv.get_lhs_offset_ex && + kernels->gemv.get_rhs_packed_offset_ex && + kernels->gemv.run_kernel_ex && + kernels->gemv.get_dst_offset; + + kernel_info * kernel = use_gemv ? &kernels->gemv : &kernels->gemm; lhs_packing_info * lhs_info = &kernels->gemm_lhs_info; - if (!kernel || !lhs_info || !lhs_info->get_offset || !lhs_info->get_packed_offset_ex || - !lhs_info->packed_size_ex || !lhs_info->pack_func_ex || + if (!kernel || !kernel->get_lhs_offset_ex || !kernel->get_rhs_packed_offset_ex || !kernel->run_kernel_ex || !kernel->get_dst_offset) { return false; } + if (!use_gemv && (!lhs_info || !lhs_info->get_offset || !lhs_info->get_packed_offset_ex || + !lhs_info->packed_size_ex || !lhs_info->pack_func_ex)) { + return false; + } + const kleidiai_weight_header * header = kleidiai_weight_header_from_ptr(src0->data); const bool has_header = kleidiai_is_weight_header_valid(header); @@ -832,16 +854,14 @@ class tensor_traits : public ggml::cpu::tensor_traits { const int nth = params->nth > 0 ? params->nth : 1; const int ith = params->ith; - const size_t k = ne00; - const size_t m = ne11; - const size_t n = ne01; - const size_t mr = kernel->get_mr(); const size_t kr = kernel->get_kr(); const size_t sr = kernel->get_sr(); - const size_t lhs_packed_size = lhs_info->packed_size_ex(m, k, 0, mr, kr, sr); - GGML_ASSERT(lhs_packed_size <= params->wsize); + const size_t lhs_packed_size = use_gemv ? 0 : lhs_info->packed_size_ex(m, k, 0, mr, kr, sr); + if (!use_gemv) { + GGML_ASSERT(lhs_packed_size <= params->wsize); + } uint8_t * lhs_packed = static_cast(params->wdata); const size_t dst_stride = dst->nb[1]; @@ -853,7 +873,7 @@ class tensor_traits : public ggml::cpu::tensor_traits { const uint8_t * lhs_batch_base = static_cast(src1->data) + batch_idx * src1->nb[2]; uint8_t * dst_batch_base = static_cast(dst->data) + batch_idx * dst->nb[2]; - { + if (!use_gemv) { const int64_t m_roundup_mr = kai_roundup((int64_t)m, (int64_t)mr); int64_t max_threads = mr ? (m_roundup_mr / (int64_t)mr) : nth; max_threads = std::max(1, max_threads); @@ -903,15 +923,17 @@ class tensor_traits : public ggml::cpu::tensor_traits { const size_t n_to_process = std::min(chunk_cols, n - n_start); if (n_to_process > 0) { - const size_t lhs_packed_offset = lhs_info->get_packed_offset_ex(0, k, 0, mr, kr, sr); + const size_t lhs_offset = use_gemv ? kernel->get_lhs_offset_ex(0, k, 0) + : lhs_info->get_packed_offset_ex(0, k, 0, mr, kr, sr); const size_t rhs_packed_offset = kernel->get_rhs_packed_offset_ex(n_start, k, 0); const size_t dst_offset = kernel->get_dst_offset(0, n_start, dst_stride); - const void * lhs_ptr = lhs_packed + lhs_packed_offset; + const void * lhs_ptr = use_gemv ? lhs_batch_base + lhs_offset + : lhs_packed + lhs_offset; const void * rhs_ptr = rhs_base + rhs_packed_offset; float * dst_ptr = reinterpret_cast(dst_batch_base + dst_offset); - kernel->run_kernel_ex(m, n_to_process, k, 0, + kernel->run_kernel_ex(m, n_to_process, k, use_gemv ? src1->nb[1] : 0, lhs_ptr, rhs_ptr, dst_ptr, From 17197474510622a3b4ea7d0909d70b606f542b96 Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Fri, 21 Aug 2026 11:33:40 +0300 Subject: [PATCH 13/38] ci : release clean-up (#27477) --- .github/workflows/make-release.yml | 12 ++++++------ .github/workflows/release.yml | 1 + scripts/make-release-desc.sh | 4 ++-- 3 files changed, 9 insertions(+), 8 deletions(-) diff --git a/.github/workflows/make-release.yml b/.github/workflows/make-release.yml index 78d7caa7ae..95e0f34755 100644 --- a/.github/workflows/make-release.yml +++ b/.github/workflows/make-release.yml @@ -62,15 +62,15 @@ jobs: GITHUB_TOKEN: ${{ github.token }} with: tag_name: ${{ steps.checks.outputs.version }} - # TODO: remove the prerelease flag once the semantic versioning workflow is ready - # ref: https://github.com/ggml-org/ggml/discussions/1579 - prerelease: true + prerelease: false + # TODO: enrich the body of the release with more information body: | - > [!NOTE] - > Semantic versioning is still work in progress. - > More info can be found in https://github.com/ggml-org/ggml/discussions/1579 + ## Overview + + New version has been released. ${{ steps.desc.outputs.nightly }} + **More info:** [dist : releases and versioning of ggml-org projects](https://github.com/ggml-org/ggml/discussions/1579) ## ${{ steps.desc.outputs.changelog_title }} diff --git a/.github/workflows/release.yml b/.github/workflows/release.yml index 61b2f5485d..a3c67604e8 100644 --- a/.github/workflows/release.yml +++ b/.github/workflows/release.yml @@ -1688,6 +1688,7 @@ jobs: GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }} with: tag_name: ${{ steps.tag.outputs.name }} + prerelease: true body: |
diff --git a/scripts/make-release-desc.sh b/scripts/make-release-desc.sh index 100f855b35..f1e4566e4a 100755 --- a/scripts/make-release-desc.sh +++ b/scripts/make-release-desc.sh @@ -52,10 +52,10 @@ PREV="$( { git tag --list; echo "${VERSION}"; } \ if [[ -n "${PREV}" ]]; then CHANGELOG="$(git log --oneline "${PREV}..${RELEASE_COMMIT}")" - CHANGELOG_TITLE="Change log since ${PREV}" + CHANGELOG_TITLE="Changelog since ${PREV}" else CHANGELOG="(no previous release tag found)" - CHANGELOG_TITLE="Change log" + CHANGELOG_TITLE="Changelog" fi # Nightly release: the b* tag pointing at the release commit (|| true: no match is not an error) From e467c2ff6174835b4079d0acf123c5ca64615088 Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Fri, 21 Aug 2026 13:20:44 +0300 Subject: [PATCH 14/38] ci : add nightly-tag.txt to make-release (#27485) As agreed in ggml discussion #1579, the official semver releases now include a nightly-tag.txt asset containing the tag of the corresponding nightly release (e.g. b10485). The Web UI assets are published to the HF bucket under the nightly tag, so this makes them discoverable for each official release. - make-release-desc.sh: expose the resolved nightly tag as a nightly_tag output - make-release.yml: create nightly-tag.txt from that tag, upload it as a release asset (skipped on dry-run), mention it in the release body and in the dry-run summary Assisted-by: pi:llama.cpp/Qwen3.8-27B --- .github/workflows/make-release.yml | 37 ++++++++++++++++++++++++++++++ scripts/make-release-desc.sh | 4 +++- 2 files changed, 40 insertions(+), 1 deletion(-) diff --git a/.github/workflows/make-release.yml b/.github/workflows/make-release.yml index 95e0f34755..451ec261f6 100644 --- a/.github/workflows/make-release.yml +++ b/.github/workflows/make-release.yml @@ -55,6 +55,20 @@ jobs: env: GITHUB_REPOSITORY: ${{ github.repository }} + - name: Create nightly-tag.txt + id: nightly_tag_file + run: | + NIGHTLY_TAG="${{ steps.desc.outputs.nightly_tag }}" + if [[ -z "${NIGHTLY_TAG}" ]]; then + echo "Warning: no nightly tag found for the release commit - nightly-tag.txt will not be created" + echo "create=false" >> "$GITHUB_OUTPUT" + exit 0 + fi + echo "${NIGHTLY_TAG}" > nightly-tag.txt + echo "create=true" >> "$GITHUB_OUTPUT" + echo "nightly-tag.txt:" + cat nightly-tag.txt + - name: Create release if: ${{ github.event.inputs.dry_run == 'false' }} uses: ggml-org/action-create-release@v1 @@ -70,18 +84,41 @@ jobs: New version has been released. ${{ steps.desc.outputs.nightly }} + + **Web UI:** the `nightly-tag.txt` asset contains the tag of the corresponding nightly release + **More info:** [dist : releases and versioning of ggml-org projects](https://github.com/ggml-org/ggml/discussions/1579) ## ${{ steps.desc.outputs.changelog_title }} ${{ steps.desc.outputs.changelog }} + - name: Upload nightly-tag.txt + if: ${{ github.event.inputs.dry_run == 'false' && steps.nightly_tag_file.outputs.create == 'true' }} + uses: actions/github-script@v8 + with: + github-token: ${{secrets.GITHUB_TOKEN}} + script: | + const fs = require('fs'); + const release_id = '${{ steps.create_release.outputs.id }}'; + console.log('uploadReleaseAsset', 'nightly-tag.txt'); + await github.rest.repos.uploadReleaseAsset({ + owner: context.repo.owner, + repo: context.repo.repo, + release_id: release_id, + name: 'nightly-tag.txt', + data: await fs.readFileSync('./nightly-tag.txt') + }); + - name: Dry run summary if: ${{ github.event.inputs.dry_run == 'true' }} run: | if [[ "${{ steps.checks.outputs.checks_passed }}" == "true" ]]; then echo "Dry run complete - all checks passed." echo "Would have created tag: ${{ steps.checks.outputs.version }}" + if [[ -n "${{ steps.desc.outputs.nightly_tag }}" ]]; then + echo "Would have uploaded nightly-tag.txt: ${{ steps.desc.outputs.nightly_tag }}" + fi else echo "::error::Dry run found release check failures. A release tag would not be created." exit 1 diff --git a/scripts/make-release-desc.sh b/scripts/make-release-desc.sh index f1e4566e4a..59aa67cba7 100755 --- a/scripts/make-release-desc.sh +++ b/scripts/make-release-desc.sh @@ -15,7 +15,8 @@ # tag exists. # # Env (when running in GitHub Actions): -# GITHUB_OUTPUT: previous_tag, changelog_title, changelog and nightly are written here +# GITHUB_OUTPUT: previous_tag, changelog_title, changelog, nightly and nightly_tag +# are written here # GITHUB_REPOSITORY: owner/repo, used to build the nightly release URL (skipped when unset) set -euo pipefail @@ -80,6 +81,7 @@ if [[ -n "${GITHUB_OUTPUT:-}" ]]; then echo "previous_tag=${PREV}" echo "changelog_title=${CHANGELOG_TITLE}" echo "nightly=${NIGHTLY}" + echo "nightly_tag=${NIGHTLY_TAG}" echo "changelog< Date: Fri, 21 Aug 2026 12:30:03 +0200 Subject: [PATCH 15/38] ui: Settings navigation cleanup (#27241) * ui : rework the settings registry into ordered raw-data sections SETTINGS_REGISTRY becomes an ordered SettingsSectionEntry[] array; the array order is the sidebar display order. Section titles, color mode options and title radio options are declared inline in their section or entry. Entries gain showInUi; MCP servers, the system-message toggle and the title LLM flag become hidden entries of their own section. Derived values (config defaults, help info, chat sections, numeric field lists, syncable parameters) are still derived here; they move to their actual consumers in follow-up commits. * ui : extract settings localStorage persistence into SettingsService Stateless load/save of the settings config and user-override keys, plus the legacy theme key migration. Business logic (default merging, mobile sendOnEnter default, applying the migrated theme) stays in the store. * ui : move the settings exit route into ROUTES SETTINGS_FALLBACK_EXIT_ROUTE is just a route, so it lives with the other routes as ROUTES.SETTINGS_EXIT. * ui : derive the syncable parameter list in the parameter sync service The syncable parameter mapping is only consumed by the sync service, so derive it there from the registry instead of exporting it from the constants file. * ui : restore isPrivate for API key masking * ui : clean up settings registry and router fetch guard Drop the per-entry section field (duplicates the parent slug and is never read) and guard the router model fetch on fields?.length so the Tools/Import-Export pages with empty fields are excluded again. Assisted-by: pi * ui : merge sampling and penalties settings into one section Assisted-by: pi --- .../settings/SettingsChat/SettingsChat.svelte | 10 +- tools/ui/src/lib/constants/index.ts | 2 +- .../ui/src/lib/constants/routes.constants.ts | 14 +- ...try.constants.ts => settings.constants.ts} | 1005 ++++++++--------- tools/ui/src/lib/services/index.ts | 10 + .../lib/services/parameter-sync.service.ts | 16 +- tools/ui/src/lib/services/settings.service.ts | 76 ++ .../src/lib/stores/settings/index.svelte.ts | 76 +- .../lib/stores/settings/referrer.svelte.ts | 4 +- tools/ui/src/lib/types/settings.d.ts | 4 +- tools/ui/src/routes/settings/+layout.svelte | 4 +- 11 files changed, 610 insertions(+), 611 deletions(-) rename tools/ui/src/lib/constants/{settings-registry.constants.ts => settings.constants.ts} (74%) create mode 100644 tools/ui/src/lib/services/settings.service.ts diff --git a/tools/ui/src/lib/components/app/settings/SettingsChat/SettingsChat.svelte b/tools/ui/src/lib/components/app/settings/SettingsChat/SettingsChat.svelte index 4233039eff..97ff30ba7a 100644 --- a/tools/ui/src/lib/components/app/settings/SettingsChat/SettingsChat.svelte +++ b/tools/ui/src/lib/components/app/settings/SettingsChat/SettingsChat.svelte @@ -15,7 +15,7 @@ NUMERIC_FIELDS, POSITIVE_INTEGER_FIELDS, SETTINGS_CHAT_SECTIONS, - SETTINGS_SECTION_TITLES + SETTINGS_SECTION_SLUGS } from '$lib/constants'; import { ColorMode } from '$lib/enums/ui.enums'; import { RouterService } from '$lib/services/router.service'; @@ -46,7 +46,7 @@ let fetchInitiated = false; $effect(() => { - if (serverStore.isRouterMode && currentSection.fields && !fetchInitiated) { + if (serverStore.isRouterMode && currentSection.fields?.length && !fetchInitiated) { fetchInitiated = true; void modelsStore @@ -148,9 +148,9 @@

{currentSection.title}

- {#if currentSection.title === SETTINGS_SECTION_TITLES.TOOLS} + {#if currentSection.slug === SETTINGS_SECTION_SLUGS.TOOLS} - {:else if currentSection.title === SETTINGS_SECTION_TITLES.IMPORT_EXPORT} + {:else if currentSection.slug === SETTINGS_SECTION_SLUGS.IMPORT_EXPORT} {:else if currentSection.fields}
@@ -161,7 +161,7 @@ onThemeChange={handleThemeChange} /> - {#if currentSection.title === SETTINGS_SECTION_TITLES.GENERAL} + {#if currentSection.slug === SETTINGS_SECTION_SLUGS.GENERAL}