#include <algorithm>
+// mul_mat(gate) + mul_mat(up) + GLU: graph shape and tensor properties only. Backend state
+// (weight layout, split buffers, DMMV) is checked by ggml_sycl_mul_mat_glu_mmvq_fused().
+static bool ggml_sycl_should_fuse_mul_mat_glu(const ggml_tensor * gate, const ggml_tensor * up,
+ const ggml_tensor * glu) {
+ // the fused epilogue implements these two; the rest fall back to the standalone GLU kernels
+ const ggml_glu_op glu_op = ggml_get_glu_op(glu);
+ if (glu_op != GGML_GLU_OP_SWIGLU && glu_op != GGML_GLU_OP_GEGLU) {
+ return false;
+ }
+
+ // the kernel always treats src[0] as the activated operand and src[1] as the multiplier
+ if (ggml_get_op_params_i32(glu, 1) /* swapped */) {
+ return false;
+ }
+
+ const ggml_tensor * wu = up->src[0];
+ 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)) {
+ 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) {
+ return false;
+ }
+
+ // one 2D reorder-layout matrix in, a plain column stride out: no broadcast or padding
+ if (!ggml_is_contiguous(wu) || !ggml_is_contiguous(wg) || !ggml_is_contiguous(act) ||
+ !ggml_is_contiguous(glu)) {
+ return false;
+ }
+ if (act->type != GGML_TYPE_F32 || glu->type != GGML_TYPE_F32) {
+ return false;
+ }
+ if (act->ne[2] != 1 || act->ne[3] != 1 || wu->ne[2] != 1 || wu->ne[3] != 1) {
+ return false;
+ }
+ // the kernel writes rows [0, wu->ne[1]) of each glu column, strided by glu->ne[0]
+ if (glu->ne[0] != wu->ne[1] || glu->ne[1] != act->ne[1]) {
+ return false;
+ }
+ // mat-vec only: one column per decoded token, up to the batch the reorder kernels cover
+ if (act->ne[1] > MMVQ_MAX_BATCH_SIZE) {
+ return false;
+ }
+
+ return true;
+}
+
bool ggml_sycl_can_fuse(const ggml_cgraph * cgraph, int node_idx, std::initializer_list<enum ggml_op> ops,
std::initializer_list<enum ggml_unary_op> unary_ops) {
#ifndef NDEBUG
return false;
}
+ // gate and up are siblings, not a chain, so ggml_can_fuse cannot express this: use the
+ // subgraph form with the GLU as the only materialised output.
+ if (ops.size() == 3 && ops.begin()[0] == GGML_OP_MUL_MAT && ops.begin()[1] == GGML_OP_MUL_MAT &&
+ ops.begin()[2] == GGML_OP_GLU) {
+ if (!ggml_can_fuse_subgraph(cgraph, node_idx, ops, { node_idx + 2 })) {
+ return false;
+ }
+
+ const ggml_tensor * glu = cgraph->nodes[node_idx + 2];
+ const ggml_tensor * gate = glu->src[0];
+ const ggml_tensor * up = glu->src[1];
+
+ // don't assume which of the two mat-muls is the gate; infer it from the GLU's operands
+ const bool ok = (gate == cgraph->nodes[node_idx] && up == cgraph->nodes[node_idx + 1]) ||
+ (gate == cgraph->nodes[node_idx + 1] && up == cgraph->nodes[node_idx]);
+ if (!ok) {
+ return false;
+ }
+
+ return ggml_sycl_should_fuse_mul_mat_glu(gate, up, glu);
+ }
+
if (!ggml_can_fuse(cgraph, node_idx, ops)) {
return false;
}
}
}
+// 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) {
+ if (!ggml_sycl_can_fuse(cgraph, node_idx, { GGML_OP_MUL_MAT, GGML_OP_MUL_MAT, GGML_OP_GLU }, {})) {
+ return false;
+ }
+
+ ggml_tensor * glu = cgraph->nodes[node_idx + 2];
+ ggml_tensor * gate = glu->src[0];
+ ggml_tensor * up = glu->src[1];
+ const ggml_tensor * wu = up->src[0];
+ const ggml_tensor * wg = gate->src[0];
+ const ggml_tensor * act = up->src[1];
+
+ // this writes glu->data directly rather than the per-device row slices that
+ // ggml_sycl_op_mul_mat() stitches back together, so it cannot serve split weights
+ if (ggml_backend_buffer_is_sycl_split(wu->buffer) || ggml_backend_buffer_is_sycl_split(wg->buffer)) {
+ return false;
+ }
+
+ // with DMMV prioritised the unfused path would not have gone through mmvq at all
+ if (g_ggml_sycl_prioritize_dmmv) {
+ return false;
+ }
+
+ // 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);
+ opt_for_reorder(&ctx, wg, act, gate, mul_mat_algo::MMVQ);
+
+ 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_g || !extra_u->optimized_feature.reorder || !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");
+
+ 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);
+
+ // one activation, quantized once and fully consumed into src1_ddq before the GEMV on this
+ // in-order queue, so glu->data aliasing the dead activation needs no memory-range check
+ 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_and_reorder_q8_1_soa>((const float *) act->data, src1_ddq, (int) ne00, (int) ne11,
+ src1_padded_cols, stream);
+
+ return ggml_sycl_mul_mat_vec_q_glu_reorder(wu->type, ggml_get_glu_op(glu), wu->data, wg->data, src1_ddq,
+ (float *) glu->data, (int) ne00, (int) wu->ne[1], (int) ne11,
+ /*stride_col_y_bytes=*/src1_padded_cols * (int) sizeof(block_q8_1) /
+ QK8_1,
+ /*stride_col_dst=*/(int) glu->ne[0], stream);
+}
__dpct_inline__ static void k_copy_src1_to_contiguous(
const char *__restrict__ src1_original, char *__restrict__ src1_contiguous,
continue;
}
+ if (node->op == GGML_OP_MUL_MAT && ggml_sycl_mul_mat_glu_mmvq_fused(*sycl_ctx, cgraph, i)) {
+ i += 2;
+ continue;
+ }
+
bool ok = ggml_sycl_compute_forward(*sycl_ctx, node);
if (!ok) {
GGML_LOG_ERROR("%s: error: op not supported %s (%s)\n", __func__, node->name, ggml_op_name(node->op));
#include "ggml.h"
#include "common.hpp"
+#include "element_wise.hpp"
#include "quants.hpp"
#include "vecdotq.hpp"
}
}
-template <typename reorder_vec_dot_q_sycl, int ncols_dst>
-static void mul_mat_vec_q_reorder_ncols(const void * __restrict__ vx, const void * __restrict__ vy,
- float * __restrict__ dst, const int ncols, const int nrows,
- const int stride_col_y_bytes, const int stride_col_dst,
- const sycl::nd_item<3> & nd_item) {
+// With has_fusion, `vgate` is a second weight matrix sharing vx's shape, stride and reorder
+// layout: one pass computes both row dot products and the epilogue writes glu(gate, up).
+template <typename reorder_vec_dot_q_sycl, int ncols_dst, bool has_fusion = false>
+static void mul_mat_vec_q_reorder_ncols(const void * __restrict__ vx, const void * __restrict__ vgate,
+ const void * __restrict__ vy, float * __restrict__ dst, const int ncols,
+ const int nrows, const int stride_col_y_bytes, const int stride_col_dst,
+ const ggml_glu_op glu_op, const sycl::nd_item<3> & nd_item) {
using block_type = ggml_sycl_reordered::block_q_t<reorder_vec_dot_q_sycl::gtype>;
using block_traits = typename block_type::traits;
const int sg_id = sg.get_group_linear_id();
const int row = workgroup_id * sg_range + sg_id;
+ // row is sub-group uniform, so this retires whole sub-groups and the collectives below
+ // stay convergent
if (row >= nrows) {
return;
}
static_assert(blocks_per_subgroup > 0);
static_assert(block_elements_per_subgroup > 0);
- float partial_sum[ncols_dst] = {0.0f};
+ float partial_sum[ncols_dst] = { 0.0f };
+ // sized 1 rather than 0 when unused: zero-length arrays are not standard C++, and the
+ // array is dead and eliminated in that case
+ [[maybe_unused]] float partial_gate[has_fusion ? ncols_dst : 1] = { 0.0f };
for (int i = sg.get_local_linear_id() / block_elements_per_subgroup; i < blocks_per_row; i += blocks_per_subgroup) {
const int ibx = row * blocks_per_row + i;
+ // the offsets depend only on the block index and the matrix shape, never on the base
+ // pointer, which is what lets vgate reuse them
const auto bx_offset = block_type::get_block_offset(ibx, nblocks);
const auto d_offset = block_type::get_d_offset(nrows, ncols, ibx);
const int iby = i * block_type::block_to_q8_1_ratio();
#pragma unroll
for (int j = 0; j < ncols_dst; ++j) {
- const char * vy_j = (const char *)vy + j * stride_col_y_bytes;
- const int8_t * q8_1_quant_ptr = (const int8_t *)vy_j + iby * QK8_1;
- const sycl::half2* q8_1_ds_ptr = (const sycl::half2 *)(vy_j + ncols + iby * sizeof(sycl::half2));
+ const char * vy_j = (const char *) vy + j * stride_col_y_bytes;
+ const int8_t * q8_1_quant_ptr = (const int8_t *) vy_j + iby * QK8_1;
+ const sycl::half2 * q8_1_ds_ptr = (const sycl::half2 *) (vy_j + ncols + iby * sizeof(sycl::half2));
partial_sum[j] += reorder_vec_dot_q_sycl()(vx, bx_offset, d_offset, q8_1_quant_ptr, q8_1_ds_ptr, iqs);
+
+ if constexpr (has_fusion) {
+ partial_gate[j] +=
+ reorder_vec_dot_q_sycl()(vgate, bx_offset, d_offset, q8_1_quant_ptr, q8_1_ds_ptr, iqs);
+ }
}
}
}
for (int j = 0; j < ncols_dst; ++j) {
float sum = sycl::reduce_over_group(nd_item.get_sub_group(), partial_sum[j], std::plus<>());
+ if constexpr (has_fusion) {
+ const float gate = sycl::reduce_over_group(nd_item.get_sub_group(), partial_gate[j], std::plus<>());
+
+ // uniform across the launch; the launcher only instantiates SWIGLU and GEGLU
+ sum *= glu_op == GGML_GLU_OP_SWIGLU ? op_silu(gate) : op_gelu(gate);
+ }
+
if (sg.leader()) {
dst[j * stride_col_dst + row] = sum;
}
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<reorder_vec_dot_q_sycl<GGML_TYPE_Q4_0>, ncols_dst>(
- vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, nd_item);
+ vx, /*vgate=*/ nullptr, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst,
+ /*glu_op=*/ GGML_GLU_OP_SWIGLU, nd_item);
});
});
}
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<reorder_vec_dot_q_sycl<GGML_TYPE_Q8_0>, ncols_dst>(
- vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, nd_item);
+ vx, /*vgate=*/ nullptr, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst,
+ /*glu_op=*/ GGML_GLU_OP_SWIGLU, nd_item);
});
});
}
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<reorder_vec_dot_q_sycl<GGML_TYPE_Q3_K>, ncols_dst>(
- vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, nd_item);
+ vx, /*vgate=*/ nullptr, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst,
+ /*glu_op=*/ GGML_GLU_OP_SWIGLU, nd_item);
});
});
}
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<reorder_vec_dot_q_sycl<GGML_TYPE_Q4_K>, ncols_dst>(
- vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, nd_item);
+ vx, /*vgate=*/ nullptr, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst,
+ /*glu_op=*/ GGML_GLU_OP_SWIGLU, nd_item);
});
});
}
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<reorder_vec_dot_q_sycl<GGML_TYPE_Q5_K>, ncols_dst>(
- vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, nd_item);
+ vx, /*vgate=*/ nullptr, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst,
+ /*glu_op=*/ GGML_GLU_OP_SWIGLU, nd_item);
});
});
}
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<reorder_vec_dot_q_sycl<GGML_TYPE_Q6_K>, ncols_dst>(
- vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, nd_item);
+ vx, /*vgate=*/ nullptr, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst,
+ /*glu_op=*/ GGML_GLU_OP_SWIGLU, nd_item);
});
});
}
return false;
}
}
+
+template <typename reorder_vec_dot_q_sycl, int ncols_dst>
+static void launch_mul_mat_vec_q_reorder_glu(const void * vx, const void * vgate, const void * vy, float * dst,
+ const int ncols, const int nrows, const int stride_col_y_bytes,
+ const int stride_col_dst, const ggml_glu_op glu_op,
+ 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<reorder_vec_dot_q_sycl, ncols_dst, /*has_fusion=*/ true>(
+ vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, glu_op,
+ nd_item);
+ });
+ });
+}
+
+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,
+ dpct::queue_ptr stream) {
+ if (src0_type != GGML_TYPE_Q4_K) {
+ return false;
+ }
+ if (glu_op != GGML_GLU_OP_SWIGLU && glu_op != GGML_GLU_OP_GEGLU) {
+ return false;
+ }
+
+ using vec_dot = reorder_vec_dot_q_sycl<GGML_TYPE_Q4_K>;
+
+ switch (ncols_dst) {
+ case 1:
+ launch_mul_mat_vec_q_reorder_glu<vec_dot, 1>(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes,
+ stride_col_dst, glu_op, stream);
+ return true;
+ case 2:
+ launch_mul_mat_vec_q_reorder_glu<vec_dot, 2>(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes,
+ stride_col_dst, glu_op, stream);
+ return true;
+ case 3:
+ launch_mul_mat_vec_q_reorder_glu<vec_dot, 3>(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes,
+ stride_col_dst, glu_op, stream);
+ return true;
+ case 4:
+ launch_mul_mat_vec_q_reorder_glu<vec_dot, 4>(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes,
+ stride_col_dst, glu_op, stream);
+ return true;
+ case 5:
+ launch_mul_mat_vec_q_reorder_glu<vec_dot, 5>(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes,
+ stride_col_dst, glu_op, stream);
+ return true;
+ case 6:
+ launch_mul_mat_vec_q_reorder_glu<vec_dot, 6>(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes,
+ stride_col_dst, glu_op, stream);
+ return true;
+ case 7:
+ launch_mul_mat_vec_q_reorder_glu<vec_dot, 7>(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes,
+ stride_col_dst, glu_op, stream);
+ return true;
+ case 8:
+ launch_mul_mat_vec_q_reorder_glu<vec_dot, 8>(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes,
+ stride_col_dst, glu_op, stream);
+ return true;
+ default:
+ return false;
+ }
+}
overview / pkg / ggml / sources / llama.cpp / commitdiff