sycl: extend MMVQ GLU fusion, add rms_norm+scale and ssm_conv+silu fusions (#28931)

* sycl : extend MMVQ GLU fusion to mixed quant types; add rms_norm+scale and ssm_conv+silu fusions

* fixing spacing issue and macro converted to template function

ggml/src/ggml-sycl/fusion.cpp

@@ -22,16 +22,22 @@ static bool ggml_sycl_should_fuse_mul_mat_glu(const ggml_tensor * gate, const gg
     const ggml_tensor * wg  = gate->src[0];
     const ggml_tensor * act = up->src[1];
 
-    // one set of block offsets and one quantized activation must serve both weights
-    if (wu->type != wg->type || !ggml_are_same_shape(wu, wg) || !ggml_are_same_stride(wu, wg)) {
+    // one activation and one output indexing must serve both weights; the block types
+    // may differ, since the plain-layout fused kernel runs each operand's own vec_dot
+    // (different types then imply different byte strides, so only the shape must agree)
+    if (!ggml_are_same_shape(wu, wg)) {
         return false;
     }
     if (act != gate->src[1]) {
         return false;
     }
 
-    // only q4_K has a fused reorder GEMV so far, and it walks whole super-blocks
-    if (wu->type != GGML_TYPE_Q4_K || wu->ne[0] % QK_K != 0) {
+    // fused GEMVs walk whole QK_K super-blocks: the reorder kernel covers same-type
+    // q4_K, the plain-layout kernel covers q5_K / iq4_xs pairs incl. mixed gate/up types
+    const bool reorder_pair = wu->type == GGML_TYPE_Q4_K && wg->type == GGML_TYPE_Q4_K;
+    const bool plain_pair   = (wu->type == GGML_TYPE_Q5_K || wu->type == GGML_TYPE_IQ4_XS) &&
+                            (wg->type == GGML_TYPE_Q5_K || wg->type == GGML_TYPE_IQ4_XS);
+    if ((!reorder_pair && !plain_pair) || wu->ne[0] % QK_K != 0) {
         return false;
     }
 
@@ -256,5 +262,19 @@ bool ggml_sycl_can_fuse(const ggml_cgraph * cgraph, int node_idx, std::initializ
         return true;
     }
 
+    if (ops.size() == 2 && ops.begin()[0] == GGML_OP_RMS_NORM && ops.begin()[1] == GGML_OP_SCALE) {
+        const ggml_tensor * rms_norm = cgraph->nodes[node_idx];
+        const ggml_tensor * scale    = cgraph->nodes[node_idx + 1];
+        GGML_ASSERT(rms_norm->src[0]->type == GGML_TYPE_F32);
+        GGML_ASSERT(rms_norm->type == GGML_TYPE_F32);
+        if (scale->src[0]->type != GGML_TYPE_F32 || scale->type != GGML_TYPE_F32) {
+            return false;
+        }
+        // the fused kernel reads/writes rows flat like the unfused pair
+        if (!ggml_is_contiguous_rows(rms_norm) || !ggml_is_contiguous_rows(scale)) {
+            return false;
+        }
+        return true;
+    }
     return false;
 }

ggml/src/ggml-sycl/ggml-sycl.cpp

@@ -4863,6 +4863,43 @@ static void ggml_sycl_mul_mat(ggml_backend_sycl_context & ctx, const ggml_tensor
     }
 }
 
+// {mul_mat(gate), mul_mat(up), GLU} over the standard (non-reorder) weight layout,
+// for quant pairs the reorder kernel does not cover (mixed gate/up types, e.g. UD-Q4_K_XL's
+// iq4_xs gate + q5_K up). Two launches replace five: one shared q8_1 quantization and one
+// dual-GEMV+GLU.
+static bool ggml_sycl_mul_mat_glu_mmvq_plain(ggml_backend_sycl_context & ctx, ggml_tensor * glu,
+                                             ggml_tensor * gate, ggml_tensor * up, const ggml_tensor * wu,
+                                             const ggml_tensor * wg, const ggml_tensor * act) {
+    // weights already migrated to the reorder layout would be misread by the plain kernel
+    const auto * extra_u = static_cast<const ggml_tensor_extra_gpu *>(wu->extra);
+    const auto * extra_g = static_cast<const ggml_tensor_extra_gpu *>(wg->extra);
+    if ((extra_u && extra_u->optimized_feature.reorder) || (extra_g && extra_g->optimized_feature.reorder)) {
+        return false;
+    }
+
+    // log the up mat-mul: glu's own srcs are the two intermediates the fusion never materialises
+    scope_op_debug_print scope_dbg_print(__func__, up, /*num_src=*/2, " : fused with gate + GLU (plain layout)");
+
+    const int64_t ne00 = wu->ne[0];
+    const int64_t ne11 = act->ne[1];
+
+    const queue_ptr stream = ctx.stream();
+    const int src1_padded_cols = GGML_PAD((int) ne00, MATRIX_ROW_PADDING);
+
+    ggml_sycl_pool_alloc<char> src1_q8_alloc(ctx.pool(),
+        (size_t) ne11 * src1_padded_cols * sizeof(block_q8_1) / QK8_1);
+    char * src1_ddq = src1_q8_alloc.get();
+
+    quantize_row_q8_1_sycl<quantize_q8_1>((const float *) act->data, src1_ddq, (int) ne00, (int) ne11,
+                                          src1_padded_cols, stream);
+
+    return ggml_sycl_mul_mat_vec_q_glu_plain(wg->type, wu->type, ggml_get_glu_op(glu), wg->data, wu->data,
+                                             src1_ddq, (float *) glu->data, (int) ne00, (int) wu->ne[1],
+                                             (int) ne11,
+                                             /*stride_col_y=*/src1_padded_cols / QK8_1,
+                                             /*stride_col_dst=*/(int) glu->ne[0], stream);
+}
+
 // Fused dense-FFN mat-vec for the {mul_mat(gate), mul_mat(up), GLU} subgraph at node_idx.
 // Returns false if it declined, in which case the caller runs the three nodes normally.
 static bool ggml_sycl_mul_mat_glu_mmvq_fused(ggml_backend_sycl_context & ctx, ggml_cgraph * cgraph, int node_idx) {
@@ -4888,6 +4925,12 @@ static bool ggml_sycl_mul_mat_glu_mmvq_fused(ggml_backend_sycl_context & ctx, gg
         return false;
     }
 
+    // quant pairs the reorder kernel cannot serve (mixed gate/up types) take the
+    // standard-layout fused path instead; q4_K keeps the reorder path below
+    if (wg->type != GGML_TYPE_Q4_K || wu->type != GGML_TYPE_Q4_K) {
+        return ggml_sycl_mul_mat_glu_mmvq_plain(ctx, glu, gate, up, wu, wg, act);
+    }
+
     // install the reorder (SoA) layout the fused kernel needs, as the unfused mmvq path would;
     // a no-op once done. after the bail checks so a declined op does not pay for it.
     opt_for_reorder(&ctx, wu, act, up, mul_mat_algo::MMVQ);
@@ -6052,6 +6095,14 @@ static void ggml_backend_sycl_graph_compute_impl(ggml_backend_sycl_context * syc
             i++;
             continue;
         }
+        // qwen35 GDN l2 norms are emitted as rms_norm + scalar scale (models.h
+        // build_gdn_l2_norm), which the rms_norm+mul fusion above cannot match
+        if (node->op == GGML_OP_RMS_NORM &&
+            ggml_sycl_can_fuse(cgraph, i, { GGML_OP_RMS_NORM, GGML_OP_SCALE }, {})) {
+            ggml_sycl_op_rms_norm_scale_fused(*sycl_ctx, node, cgraph->nodes[i + 1]);
+            i++;
+            continue;
+        }
         if (node->op == GGML_OP_ADD &&
             ggml_sycl_can_fuse(cgraph, i, { GGML_OP_ADD, GGML_OP_ADD }, {})) {
             ggml_sycl_op_add_add_fused(*sycl_ctx, node, cgraph->nodes[i + 1]);

ggml/src/ggml-sycl/mmvq.cpp

@@ -3051,6 +3051,191 @@ static void launch_mul_mat_vec_q_reorder_glu(const void * vx, const void * vgate
     launch_mul_mat_vec_q_reorder_glu_impl<reorder_vec_dot_q_sycl, ncols_dst, rows_per_sg>(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, glu_op, stream);
 }
 
+// ---------------------------------------------------------------------------
+// Fused dense-FFN GEMV + GLU over the standard (non-reorder) weight layout.
+//
+// Unlike the reorder variant below, the two weights may carry different block
+// types (e.g. an unsloth UD mix with an iq4_xs gate and a q5_K up), as long as
+// both quantize in QK_K-sized super-blocks so that one q8_1 activation
+// quantization serves both dots. Per-operand accumulation order matches
+// mul_mat_vec_q exactly, so results are bit-identical to running the three
+// nodes separately.
+// ---------------------------------------------------------------------------
+template <int qi_g, typename block_g_t, int vdr_g, vec_dot_q_sycl_t vec_dot_g,
+          int qi_u, typename block_u_t, int vdr_u, vec_dot_q_sycl_t vec_dot_u, int ncols_dst>
+static void mul_mat_vec_q_glu(const void * __restrict__ vxg, const void * __restrict__ vxu,
+                              const void * __restrict__ vy, float * __restrict__ dst, const int ncols,
+                              const int nrows, const int stride_col_y, const int stride_col_dst,
+                              const ggml_glu_op glu_op, const sycl::nd_item<3> & item_ct1) {
+    static_assert(QK_K % QK8_1 == 0);
+    const int row = item_ct1.get_group(2) * item_ct1.get_local_range(1) + item_ct1.get_local_id(1);
+    if (row >= nrows) {
+        return;
+    }
+    const int     blocks_per_row  = ncols / QK_K;
+    constexpr int blocks_per_warp_g = (vdr_g * WARP_SIZE + qi_g - 1) / qi_g;
+    constexpr int blocks_per_warp_u = (vdr_u * WARP_SIZE + qi_u - 1) / qi_u;
+    // one partial sum per output column, per operand
+    float tmpg[ncols_dst] = {0.0f};
+    float tmpu[ncols_dst] = {0.0f};
+    const block_g_t  * xg = (const block_g_t *) vxg;
+    const block_u_t  * xu = (const block_u_t *) vxu;
+    const block_q8_1 * y  = (const block_q8_1 *) vy;
+    for (int i = item_ct1.get_local_id(2) / (qi_g / vdr_g); i < blocks_per_row; i += blocks_per_warp_g) {
+        const int ibx = row * blocks_per_row + i;
+        const int iby = i * (QK_K / QK8_1);
+        for (size_t elem = 0; elem < qi_g / vdr_g; elem += WARP_SIZE) {
+            const int iqs = elem + vdr_g * (item_ct1.get_local_id(2) % (qi_g / vdr_g));
+#pragma unroll
+            for (int j = 0; j < ncols_dst; ++j) {
+                tmpg[j] += vec_dot_g(&xg[ibx], &y[j * stride_col_y + iby], iqs);
+            }
+        }
+    }
+    for (int i = item_ct1.get_local_id(2) / (qi_u / vdr_u); i < blocks_per_row; i += blocks_per_warp_u) {
+        const int ibx = row * blocks_per_row + i;
+        const int iby = i * (QK_K / QK8_1);
+        for (size_t elem = 0; elem < qi_u / vdr_u; elem += WARP_SIZE) {
+            const int iqs = elem + vdr_u * (item_ct1.get_local_id(2) % (qi_u / vdr_u));
+#pragma unroll
+            for (int j = 0; j < ncols_dst; ++j) {
+                tmpu[j] += vec_dot_u(&xu[ibx], &y[j * stride_col_y + iby], iqs);
+            }
+        }
+    }
+    // sum up partial sums and write back the activated product
+#pragma unroll
+    for (int j = 0; j < ncols_dst; ++j) {
+#pragma unroll
+        for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) {
+            tmpg[j] += dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmpg[j], mask);
+            tmpu[j] += dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmpu[j], mask);
+        }
+    }
+    if (item_ct1.get_local_id(2) == 0) {
+#pragma unroll
+        for (int j = 0; j < ncols_dst; ++j) {
+            // uniform across the launch; the dispatcher only accepts SWIGLU and GEGLU
+            const float gate = glu_op == GGML_GLU_OP_SWIGLU ? op_silu(tmpg[j]) : op_gelu(tmpg[j]);
+            dst[j * stride_col_dst + row] = gate * tmpu[j];
+        }
+    }
+}
+
+template <int qi_g, typename block_g_t, int vdr_g, vec_dot_q_sycl_t vec_dot_g,
+          int qi_u, typename block_u_t, int vdr_u, vec_dot_q_sycl_t vec_dot_u, int ncols_dst>
+static void launch_mul_mat_vec_q_glu(const void * vxg, const void * vxu, const void * vy, float * dst,
+                                     const int ncols, const int nrows, const int stride_col_y,
+                                     const int stride_col_dst, const ggml_glu_op glu_op,
+                                     dpct::queue_ptr stream) {
+    GGML_ASSERT(ncols % QK_K == 0);
+    const int            block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
+    const sycl::range<3> block_nums(1, 1, block_num_y);
+    const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, 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_glu<qi_g, block_g_t, vdr_g, vec_dot_g,
+                                               qi_u, block_u_t, vdr_u, vec_dot_u, ncols_dst>(
+                                 vxg, vxu, vy, dst, ncols, nrows, stride_col_y, stride_col_dst,
+                                 glu_op, nd_item);
+                         });
+    });
+}
+
+// Dispatch the plain-layout fused GLU GEMV over the activation batch: ncols_dst
+// selects the kernel's per-column template parameter. Returns false when the
+// batch exceeds the instantiated range; the caller falls back to unfused nodes.
+template <int qi_g, typename block_g_t, int vdr_g, vec_dot_q_sycl_t vec_dot_g,
+          int qi_u, typename block_u_t, int vdr_u, vec_dot_q_sycl_t vec_dot_u>
+static bool dispatch_mul_mat_vec_q_glu_plain(const void * vgate, const void * vup, const void * vy,
+                                             float * dst, const int ncols, const int nrows,
+                                             const int stride_col_y, const int stride_col_dst,
+                                             const ggml_glu_op glu_op, dpct::queue_ptr stream,
+                                             const int ncols_dst) {
+    switch (ncols_dst) {
+        case 1:
+            launch_mul_mat_vec_q_glu<qi_g, block_g_t, vdr_g, vec_dot_g,
+                                     qi_u, block_u_t, vdr_u, vec_dot_u, 1>(
+                vgate, vup, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, glu_op, stream);
+            return true;
+        case 2:
+            launch_mul_mat_vec_q_glu<qi_g, block_g_t, vdr_g, vec_dot_g,
+                                     qi_u, block_u_t, vdr_u, vec_dot_u, 2>(
+                vgate, vup, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, glu_op, stream);
+            return true;
+        case 3:
+            launch_mul_mat_vec_q_glu<qi_g, block_g_t, vdr_g, vec_dot_g,
+                                     qi_u, block_u_t, vdr_u, vec_dot_u, 3>(
+                vgate, vup, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, glu_op, stream);
+            return true;
+        case 4:
+            launch_mul_mat_vec_q_glu<qi_g, block_g_t, vdr_g, vec_dot_g,
+                                     qi_u, block_u_t, vdr_u, vec_dot_u, 4>(
+                vgate, vup, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, glu_op, stream);
+            return true;
+        case 5:
+            launch_mul_mat_vec_q_glu<qi_g, block_g_t, vdr_g, vec_dot_g,
+                                     qi_u, block_u_t, vdr_u, vec_dot_u, 5>(
+                vgate, vup, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, glu_op, stream);
+            return true;
+        case 6:
+            launch_mul_mat_vec_q_glu<qi_g, block_g_t, vdr_g, vec_dot_g,
+                                     qi_u, block_u_t, vdr_u, vec_dot_u, 6>(
+                vgate, vup, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, glu_op, stream);
+            return true;
+        case 7:
+            launch_mul_mat_vec_q_glu<qi_g, block_g_t, vdr_g, vec_dot_g,
+                                     qi_u, block_u_t, vdr_u, vec_dot_u, 7>(
+                vgate, vup, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, glu_op, stream);
+            return true;
+        case 8:
+            launch_mul_mat_vec_q_glu<qi_g, block_g_t, vdr_g, vec_dot_g,
+                                     qi_u, block_u_t, vdr_u, vec_dot_u, 8>(
+                vgate, vup, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, glu_op, stream);
+            return true;
+        default:
+            return false;
+    }
+}
+
+// Fused dense-FFN GEMV + GLU for weight pairs the reorder kernel does not cover.
+// vgate/vup must be in the standard block layout; vy must be quantized with plain
+// quantize_q8_1 (padded rows). stride_col_y is in block_q8_1 units.
+// Returns false if the type pair or batch is unhandled; caller should fall back.
+bool ggml_sycl_mul_mat_vec_q_glu_plain(enum ggml_type gate_type, enum ggml_type up_type,
+                                       enum ggml_glu_op glu_op, const void * vgate, const void * vup,
+                                       const void * vy, float * dst, int ncols, int nrows, int ncols_dst,
+                                       int stride_col_y, int stride_col_dst, dpct::queue_ptr stream) {
+    if (glu_op != GGML_GLU_OP_SWIGLU && glu_op != GGML_GLU_OP_GEGLU) {
+        return false;
+    }
+    if (ncols % QK_K != 0) {
+        return false;
+    }
+    if (gate_type == GGML_TYPE_Q5_K && up_type == GGML_TYPE_Q5_K) {
+        return dispatch_mul_mat_vec_q_glu_plain<QI5_K, block_q5_K, VDR_Q5_K_Q8_1_MMVQ, vec_dot_q5_K_q8_1,
+                                                QI5_K, block_q5_K, VDR_Q5_K_Q8_1_MMVQ, vec_dot_q5_K_q8_1>(
+            vgate, vup, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, glu_op, stream, ncols_dst);
+    }
+    if (gate_type == GGML_TYPE_IQ4_XS && up_type == GGML_TYPE_IQ4_XS) {
+        return dispatch_mul_mat_vec_q_glu_plain<QI4_XS / 4, block_iq4_xs, VDR_IQ4_XS_Q8_1_MMVQ, vec_dot_iq4_xs_q8_1,
+                                                QI4_XS / 4, block_iq4_xs, VDR_IQ4_XS_Q8_1_MMVQ, vec_dot_iq4_xs_q8_1>(
+            vgate, vup, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, glu_op, stream, ncols_dst);
+    }
+    if (gate_type == GGML_TYPE_IQ4_XS && up_type == GGML_TYPE_Q5_K) {
+        return dispatch_mul_mat_vec_q_glu_plain<QI4_XS / 4, block_iq4_xs, VDR_IQ4_XS_Q8_1_MMVQ, vec_dot_iq4_xs_q8_1,
+                                                QI5_K, block_q5_K, VDR_Q5_K_Q8_1_MMVQ, vec_dot_q5_K_q8_1>(
+            vgate, vup, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, glu_op, stream, ncols_dst);
+    }
+    if (gate_type == GGML_TYPE_Q5_K && up_type == GGML_TYPE_IQ4_XS) {
+        return dispatch_mul_mat_vec_q_glu_plain<QI5_K, block_q5_K, VDR_Q5_K_Q8_1_MMVQ, vec_dot_q5_K_q8_1,
+                                                QI4_XS / 4, block_iq4_xs, VDR_IQ4_XS_Q8_1_MMVQ, vec_dot_iq4_xs_q8_1>(
+            vgate, vup, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, glu_op, stream, ncols_dst);
+    }
+    return false;
+}
+
 bool ggml_sycl_mul_mat_vec_q_glu_reorder(enum ggml_type src0_type, enum ggml_glu_op glu_op, const void * vx,
                                          const void * vgate, const void * vy, float * dst, int ncols, int nrows,
                                          int ncols_dst, int stride_col_y_bytes, int stride_col_dst,

ggml/src/ggml-sycl/mmvq.hpp

@@ -73,4 +73,24 @@ bool ggml_sycl_mul_mat_vec_q_glu_reorder(
     int                stride_col_dst,       // floats between output columns in dst
     dpct::queue_ptr    stream);
 
+
+// Fused dense-FFN GEMV + GLU over the standard (non-reorder) layout; the gate and up
+// weights may carry different block types (q5_K / iq4_xs, mixed included).
+// vy: src1 quantized with plain quantize_q8_1 (padded rows). stride_col_y is in
+// block_q8_1 units. Returns false if the pair or batch is unhandled; caller falls back.
+bool ggml_sycl_mul_mat_vec_q_glu_plain(
+    enum ggml_type     gate_type,
+    enum ggml_type     up_type,
+    enum ggml_glu_op   glu_op,
+    const void *       vgate,
+    const void *       vup,
+    const void *       vy,
+    float *            dst,
+    int                ncols,                // K, shared by both weights
+    int                nrows,                // output rows, i.e. weight ne[1]
+    int                ncols_dst,            // activation columns, 1..MMVQ_MAX_BATCH_SIZE
+    int                stride_col_y,         // block_q8_1 units between activation columns
+    int                stride_col_dst,       // floats between output columns in dst
+    dpct::queue_ptr    stream);
+
 #endif // GGML_SYCL_MMVQ_HPP

ggml/src/ggml-sycl/norm.cpp

@@ -152,7 +152,8 @@ 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 float* add = nullptr, const int64_t add_stride_row = 0, const int64_t add_stride_channel = 0,
-    const int64_t add_stride_sample = 0, const int add_nrows = 0, const int add_nchannels = 0, const int add_nsamples = 0) {
+    const int64_t add_stride_sample = 0, const int add_nrows = 0, const int add_nchannels = 0, const int add_nsamples = 0,
+    const float scale_mul = 1.0f) {
 
     static_assert(!do_add || do_multiply, "fusing add is not supported without multiplying");
 
@@ -221,7 +222,11 @@ static void rms_norm_f32(const float* x, float* dst, const int ncols,
         } else if constexpr (do_multiply) {
             dst[col * dst_stride_col] = scale * x[col * src_stride_col] * mul[col];
         } else {
-            dst[col * dst_stride_col] = scale * x[col * src_stride_col];
+            // folded epilogue of a fused GGML_OP_SCALE consumer (qwen35 GDN l2 norms);
+            // the explicit temporary keeps the float evaluation order identical to
+            // running rms_norm and scale as two separate kernels
+            const float v = scale * x[col * src_stride_col];
+            dst[col * dst_stride_col] = v * scale_mul;
         }
     }
 }
@@ -394,6 +399,52 @@ static void rms_norm_f32_sycl(const float* x, float* dst, const int ncols, const
     }
 }
 
+static void rms_norm_scale_f32_sycl(const float* x, float* dst, const int ncols, const int nrows,
+    const int nchannels, const int nsamples,
+    const int64_t src_stride_col, const int64_t src_stride_row, const int64_t src_stride_channel, 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 float scale_mul, queue_ptr stream, int device) {
+    const sycl::range<3> global_dims(nsamples, nchannels, nrows);
+    if (ncols < 1024) {
+        const sycl::range<3> block_dims(1, 1, WARP_SIZE);
+        stream->submit([&](sycl::handler& cgh) {
+            cgh.parallel_for(
+                sycl::nd_range<3>(global_dims * block_dims, block_dims),
+                [=](sycl::nd_item<3> item_ct1)
+                [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
+                    rms_norm_f32(x, dst, ncols,
+                        src_stride_col, src_stride_row, src_stride_channel, src_stride_sample,
+                        dst_stride_col, dst_stride_row, dst_stride_channel, dst_stride_sample,
+                        eps, item_ct1, nullptr, WARP_SIZE,
+                        nullptr, 0, 0, 0, 0, 0, 0,
+                        nullptr, 0, 0, 0, 0, 0, 0,
+                        scale_mul);
+                });
+            });
+    }
+    else {
+        const int work_group_size = ggml_sycl_info().max_work_group_sizes[device];
+        assert(work_group_size % (WARP_SIZE * WARP_SIZE) == 0);
+        const sycl::range<3> block_dims(1, 1, work_group_size);
+        stream->submit([&](sycl::handler& cgh) {
+            sycl::local_accessor<float, 1> s_sum_acc_ct1(sycl::range<1>(work_group_size / WARP_SIZE),
+                cgh);
+            cgh.parallel_for(
+                sycl::nd_range<3>(global_dims * block_dims, block_dims),
+                [=](sycl::nd_item<3> item_ct1)
+                [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
+                    rms_norm_f32(x, dst, ncols,
+                        src_stride_col, src_stride_row, src_stride_channel, src_stride_sample,
+                        dst_stride_col, dst_stride_row, dst_stride_channel, dst_stride_sample,
+                        eps, item_ct1, get_pointer(s_sum_acc_ct1), work_group_size,
+                        nullptr, 0, 0, 0, 0, 0, 0,
+                        nullptr, 0, 0, 0, 0, 0, 0,
+                        scale_mul);
+                });
+            });
+    }
+}
+
 static void rms_norm_mul_f32_sycl(const float* x, const float* mul, float* dst, const int ncols, const int nrows,
         const int nchannels, const int nsamples,
         const int64_t src_stride_col, const int64_t src_stride_row, const int64_t src_stride_channel, const int64_t src_stride_sample,
@@ -682,6 +733,46 @@ void ggml_sycl_op_rms_norm(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
         ss0, ss1, ss2, ss3, ds0, ds1, ds2, ds3, eps, main_stream, ctx.device);
 }
 
+// Fused rms_norm + scale (the qwen35 GDN l2-norm pair build_gdn_l2_norm emits):
+// the scale factor is a host scalar in the GGML_OP_SCALE node's op_params, so
+// unlike the mul variants there is no second device tensor to wire up.
+void ggml_sycl_op_rms_norm_scale_fused(ggml_backend_sycl_context & ctx, ggml_tensor * dst,
+                                       ggml_tensor * scale_tensor) {
+    const ggml_tensor * src0 = dst->src[0];
+    GGML_ASSERT(src0->type == GGML_TYPE_F32);
+    GGML_ASSERT(dst->type == GGML_TYPE_F32);
+    GGML_ASSERT(scale_tensor->type == GGML_TYPE_F32);
+
+    dpct::queue_ptr main_stream = ctx.stream();
+    SYCL_CHECK(ggml_sycl_set_device(ctx.device));
+
+    const float * src0_dd = static_cast<const float *>(src0->data);
+    float *       dst_dd  = static_cast<float *>(scale_tensor->data);
+
+    float eps;
+    memcpy(&eps, dst->op_params, sizeof(float));
+    float scale_mul;
+    memcpy(&scale_mul, scale_tensor->op_params, sizeof(float));
+    GGML_ASSERT(scale_mul >= 0.0f);
+
+    GGML_TENSOR_UNARY_OP_LOCALS
+    const size_t ts0 = ggml_type_size(src0->type);
+    const size_t tdst = ggml_type_size(scale_tensor->type);
+    GGML_ASSERT(nb00 % ts0 == 0 && nb01 % ts0 == 0 && nb02 % ts0 == 0 && nb03 % ts0 == 0);
+    GGML_ASSERT(scale_tensor->nb[0] % tdst == 0 && scale_tensor->nb[1] % tdst == 0 &&
+                scale_tensor->nb[2] % tdst == 0 && scale_tensor->nb[3] % tdst == 0);
+    const int64_t ss0 = nb00 / ts0;
+    const int64_t ss1 = nb01 / ts0;
+    const int64_t ss2 = nb02 / ts0;
+    const int64_t ss3 = nb03 / ts0;
+    const int64_t ds0 = scale_tensor->nb[0] / tdst;
+    const int64_t ds1 = scale_tensor->nb[1] / tdst;
+    const int64_t ds2 = scale_tensor->nb[2] / tdst;
+    const int64_t ds3 = scale_tensor->nb[3] / tdst;
+    rms_norm_scale_f32_sycl(src0_dd, dst_dd, ne00, ne01, ne02, ne03,
+        ss0, ss1, ss2, ss3, ds0, ds1, ds2, ds3, eps, scale_mul, main_stream, ctx.device);
+}
+
 void ggml_sycl_op_rms_norm_fused(ggml_backend_sycl_context & ctx, ggml_tensor * dst, ggml_tensor * mul_tensor) {
     const ggml_tensor * rms_norm_src = dst->src[0];
     float eps = 0.0f;

ggml/src/ggml-sycl/norm.hpp

@@ -21,6 +21,8 @@ void ggml_sycl_op_rms_norm(ggml_backend_sycl_context& ctx, ggml_tensor* dst);
 
 void ggml_sycl_op_rms_norm_fused(ggml_backend_sycl_context& ctx, ggml_tensor* dst, ggml_tensor* mul);
 
+void ggml_sycl_op_rms_norm_scale_fused(ggml_backend_sycl_context& ctx, ggml_tensor* dst, ggml_tensor* scale_tensor);
+
 void ggml_sycl_op_rms_norm_fused_add(ggml_backend_sycl_context& ctx, ggml_tensor* dst, ggml_tensor* mul_tensor, ggml_tensor* add_tensor);
 
 void ggml_sycl_op_rms_norm_back(ggml_backend_sycl_context& ctx, ggml_tensor* dst);