}
}
+static void dequantize_mul_mat_vec_q5_k_reorder(const void *__restrict__ vx,
+ const float *__restrict__ yy,
+ float *__restrict__ dst,
+ const int ncols, int nrows,
+ const sycl::nd_item<3> &item_ct1) {
+
+ const int row = item_ct1.get_group(2);
+ const int num_blocks_per_row = ncols / QK_K;
+ const int ib0 = row*num_blocks_per_row;
+
+ // SOA base pointers for the reordered layout:
+ // [qs: nb * QK_K/2] [qh: nb * QK_K/8] [scales: nb * K_SCALE_SIZE] [dm: nb * sizeof(half2)]
+ const int nb = nrows * num_blocks_per_row;
+ const uint8_t * qs_base = (const uint8_t *)vx;
+ const uint8_t * qh_base = qs_base + (size_t)nb * (QK_K / 2);
+ const uint8_t * scales_base = qh_base + (size_t)nb * (QK_K / 8);
+ const sycl::half2 * dm_base = (const sycl::half2 *)(scales_base + (size_t)nb * K_SCALE_SIZE);
+
+ float tmp = 0; // partial sum for thread in warp
+
+#if QK_K == 256
+ const uint16_t kmask1 = 0x3f3f;
+ const uint16_t kmask2 = 0x0f0f;
+ const uint16_t kmask3 = 0xc0c0;
+
+ const int tid = item_ct1.get_local_id(2) / 2; // 0...15
+ const int ix = item_ct1.get_local_id(2) % 2;
+
+ const int il = tid/4; // 0...3
+ const int ir = tid - 4*il;// 0...3
+ const int n = 2;
+
+ const int im = il/2; // 0 or 1. 0 computes 0,32 + 128,160, 1 computes 64,96 + 192,224
+ const int in = il%2;
+
+ const int l0 = n*(2*ir + in);
+ const int q_offset = 32*im + l0;
+ const int y_offset = 64*im + l0;
+
+ const uint8_t hm1 = 1 << (2*im);
+ const uint8_t hm2 = hm1 << 4;
+
+ uint16_t aux[4];
+ const uint8_t * sc = (const uint8_t *)aux;
+
+ uint16_t q16[8];
+ const uint8_t * q4 = (const uint8_t *)q16;
+
+ for (int i = ix; i < num_blocks_per_row; i += 2) {
+ const int bi = ib0 + i;
+
+ const uint8_t * ql1 = qs_base + bi * (QK_K / 2) + q_offset;
+ const uint8_t * qh = qh_base + bi * (QK_K / 8) + l0;
+ const float * y1 = yy + i*QK_K + y_offset;
+ const float * y2 = y1 + 128;
+
+ const sycl::half2 dm_val = dm_base[bi];
+ const float dall = dm_val[0];
+ const float dmin = dm_val[1];
+
+ const uint16_t * a = (const uint16_t *)(scales_base + bi * K_SCALE_SIZE);
+ aux[0] = a[im+0] & kmask1;
+ aux[1] = a[im+2] & kmask1;
+ aux[2] = ((a[im+4] >> 0) & kmask2) | ((a[im+0] & kmask3) >> 2);
+ aux[3] = ((a[im+4] >> 4) & kmask2) | ((a[im+2] & kmask3) >> 2);
+
+ sycl::float4 sum = {0.f, 0.f, 0.f, 0.f};
+ float smin = 0;
+ const uint16_t * q1 = (const uint16_t *)ql1;
+ const uint16_t * q2 = q1 + 32;
+ q16[0] = q1[0] & 0x0f0f;
+ q16[1] = q1[8] & 0x0f0f;
+ q16[2] = (q1[0] >> 4) & 0x0f0f;
+ q16[3] = (q1[8] >> 4) & 0x0f0f;
+ q16[4] = q2[0] & 0x0f0f;
+ q16[5] = q2[8] & 0x0f0f;
+ q16[6] = (q2[0] >> 4) & 0x0f0f;
+ q16[7] = (q2[8] >> 4) & 0x0f0f;
+ for (int l = 0; l < n; ++l) {
+ sum.x() +=
+ y1[l + 0] * (q4[l + 0] + (qh[l + 0] & (hm1 << 0) ? 16 : 0)) +
+ y1[l + 16] * (q4[l + 2] + (qh[l + 16] & (hm1 << 0) ? 16 : 0));
+ sum.y() +=
+ y1[l + 32] * (q4[l + 4] + (qh[l + 0] & (hm1 << 1) ? 16 : 0)) +
+ y1[l + 48] * (q4[l + 6] + (qh[l + 16] & (hm1 << 1) ? 16 : 0));
+ sum.z() +=
+ y2[l + 0] * (q4[l + 8] + (qh[l + 0] & (hm2 << 0) ? 16 : 0)) +
+ y2[l + 16] * (q4[l + 10] + (qh[l + 16] & (hm2 << 0) ? 16 : 0));
+ sum.w() +=
+ y2[l + 32] * (q4[l + 12] + (qh[l + 0] & (hm2 << 1) ? 16 : 0)) +
+ y2[l + 48] * (q4[l + 14] + (qh[l + 16] & (hm2 << 1) ? 16 : 0));
+ smin += (y1[l] + y1[l+16]) * sc[2] + (y1[l+32] + y1[l+48]) * sc[3]
+ + (y2[l] + y2[l+16]) * sc[6] + (y2[l+32] + y2[l+48]) * sc[7];
+ }
+ tmp += dall * (sum.x() * sc[0] + sum.y() * sc[1] + sum.z() * sc[4] +
+ sum.w() * sc[5]) -
+ dmin * smin;
+ }
+#else
+ // The reordered Q5_K layout is only produced for QK_K == 256.
+#endif
+
+ // sum up partial sums and write back result
+#pragma unroll
+ for (int mask = QK_WARP_SIZE / 2; mask > 0; mask >>= 1) {
+ tmp +=
+ dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask);
+ }
+
+ if (item_ct1.get_local_id(2) == 0) {
+ dst[row] = tmp;
+ }
+}
+
static void dequantize_mul_mat_vec_q6_k(const void * __restrict__ vx, const float * __restrict__ yy, float * __restrict__ dst, const int ncols, int nrows,
const sycl::nd_item<3> &item_ct1) {
});
}
+static void dequantize_mul_mat_vec_q5_K_sycl_reorder(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 sycl::range<3> block_dims(1, 1, QK_WARP_SIZE);
+ stream->parallel_for(
+ sycl::nd_range<3>(sycl::range<3>(1, 1, nrows) * block_dims, block_dims),
+ [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] {
+ dequantize_mul_mat_vec_q5_k_reorder(vx, y, dst, ncols, nrows, item_ct1);
+ });
+}
+
static void dequantize_mul_mat_vec_q6_K_sycl_reorder(const void *vx, const float *y,
float *dst, const int ncols,
const int nrows,
}
break;
case GGML_TYPE_Q5_K:
- dequantize_mul_mat_vec_q5_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ 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);
+ } else {
+ dequantize_mul_mat_vec_q5_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ }
break;
case GGML_TYPE_Q6_K:
if ((ggml_tensor_extra_gpu *) dst->src[0]->extra &&
return true;
}
+// Reorder each expert slice into a self-contained SoA layout.
+static bool reorder_qw_q4_k_moe(uint8_t * data_device, size_t expert_bytes, int64_t n_expert, dpct::queue_ptr stream) {
+ GGML_ASSERT(expert_bytes % sizeof(block_q4_K) == 0);
+ const int blocks_per_expert = (int) (expert_bytes / sizeof(block_q4_K));
+ const size_t total_bytes = expert_bytes * (size_t) n_expert;
+
+ sycl_reorder_temp_buffer tmp(stream, total_bytes);
+ if (!tmp) {
+ GGML_LOG_WARN("%s: failed to allocate %zu bytes for reorder temp buffer, skipping reorder\n", __func__, total_bytes);
+ return false;
+ }
+ uint8_t * tmp_buf = static_cast<uint8_t *>(tmp.ptr);
+
+ sycl::event copy_event;
+ SYCL_CHECK(CHECK_TRY_ERROR(copy_event = stream->memcpy(tmp_buf, data_device, total_bytes)));
+ if (!g_ggml_sycl_use_async_mem_op) {
+ copy_event.wait();
+ }
+
+ const int total_blocks = blocks_per_expert * (int) n_expert;
+ auto reorder_event = stream->parallel_for(total_blocks, [=](auto gb_) {
+ const int gb = gb_;
+ const int e = gb / blocks_per_expert;
+ const int ib = gb % blocks_per_expert;
+ const block_q4_K * x = (const block_q4_K *) (tmp_buf + (size_t) e * expert_bytes);
+ uint8_t * base = data_device + (size_t) e * expert_bytes;
+
+ auto * qs_ptr = base;
+ auto * scales_ptr = qs_ptr + QK_K / 2 * blocks_per_expert;
+ auto * dm_ptr = (sycl::half2 *) (scales_ptr + K_SCALE_SIZE * blocks_per_expert);
+
+ for (int j = 0; j < QK_K / 2; ++j) {
+ qs_ptr[ib * (QK_K / 2) + j] = x[ib].qs[j];
+ }
+ for (int j = 0; j < K_SCALE_SIZE; ++j) {
+ scales_ptr[ib * K_SCALE_SIZE + j] = x[ib].scales[j];
+ }
+ dm_ptr[ib] = x[ib].dm;
+ });
+ if (!g_ggml_sycl_use_async_mem_op) {
+ reorder_event.wait_and_throw();
+ }
+ return true;
+}
+
+// Reorder each Q5_K expert slice into [qs][qh][scales][dm].
+static bool reorder_qw_q5_k_moe(uint8_t * data_device, size_t expert_bytes, int64_t n_expert, dpct::queue_ptr stream) {
+ GGML_ASSERT(expert_bytes % sizeof(block_q5_K) == 0);
+ const int blocks_per_expert = (int) (expert_bytes / sizeof(block_q5_K));
+ const size_t total_bytes = expert_bytes * (size_t) n_expert;
+
+ sycl_reorder_temp_buffer tmp(stream, total_bytes);
+ if (!tmp) {
+ GGML_LOG_WARN("%s: failed to allocate %zu bytes for reorder temp buffer, skipping reorder\n", __func__, total_bytes);
+ return false;
+ }
+ uint8_t * tmp_buf = static_cast<uint8_t *>(tmp.ptr);
+
+ sycl::event copy_event;
+ SYCL_CHECK(CHECK_TRY_ERROR(copy_event = stream->memcpy(tmp_buf, data_device, total_bytes)));
+ if (!g_ggml_sycl_use_async_mem_op) {
+ copy_event.wait();
+ }
+
+ const int total_blocks = blocks_per_expert * (int) n_expert;
+ auto reorder_event = stream->parallel_for(total_blocks, [=](auto gb_) {
+ const int gb = gb_;
+ const int e = gb / blocks_per_expert;
+ const int ib = gb % blocks_per_expert;
+ const block_q5_K * x = (const block_q5_K *) (tmp_buf + (size_t) e * expert_bytes);
+ uint8_t * base = data_device + (size_t) e * expert_bytes;
+
+ auto * qs_ptr = base;
+ auto * qh_ptr = qs_ptr + (QK_K / 2) * blocks_per_expert;
+ auto * scales_ptr = qh_ptr + (QK_K / 8) * blocks_per_expert;
+ auto * dm_ptr = (sycl::half2 *) (scales_ptr + K_SCALE_SIZE * blocks_per_expert);
+
+ for (int j = 0; j < QK_K / 2; ++j) {
+ qs_ptr[ib * (QK_K / 2) + j] = x[ib].qs[j];
+ }
+ for (int j = 0; j < QK_K / 8; ++j) {
+ qh_ptr[ib * (QK_K / 8) + j] = x[ib].qh[j];
+ }
+ for (int j = 0; j < K_SCALE_SIZE; ++j) {
+ scales_ptr[ib * K_SCALE_SIZE + j] = x[ib].scales[j];
+ }
+ dm_ptr[ib] = x[ib].dm;
+ });
+ if (!g_ggml_sycl_use_async_mem_op) {
+ reorder_event.wait_and_throw();
+ }
+ return true;
+}
+
+// Reorder each Q6_K expert slice into [ql][qh][scales][d].
+static bool reorder_qw_q6_k_moe(uint8_t * data_device, size_t expert_bytes, int64_t n_expert, dpct::queue_ptr stream) {
+ GGML_ASSERT(expert_bytes % sizeof(block_q6_K) == 0);
+ const int blocks_per_expert = (int) (expert_bytes / sizeof(block_q6_K));
+ const size_t total_bytes = expert_bytes * (size_t) n_expert;
+
+ sycl_reorder_temp_buffer tmp(stream, total_bytes);
+ if (!tmp) {
+ GGML_LOG_WARN("%s: failed to allocate %zu bytes for reorder temp buffer, skipping reorder\n", __func__, total_bytes);
+ return false;
+ }
+ uint8_t * tmp_buf = static_cast<uint8_t *>(tmp.ptr);
+
+ sycl::event copy_event;
+ SYCL_CHECK(CHECK_TRY_ERROR(copy_event = stream->memcpy(tmp_buf, data_device, total_bytes)));
+ if (!g_ggml_sycl_use_async_mem_op) {
+ copy_event.wait();
+ }
+
+ const int total_blocks = blocks_per_expert * (int) n_expert;
+ auto reorder_event = stream->parallel_for(total_blocks, [=](auto gb_) {
+ const int gb = gb_;
+ const int e = gb / blocks_per_expert;
+ const int ib = gb % blocks_per_expert;
+ const block_q6_K * x = (const block_q6_K *) (tmp_buf + (size_t) e * expert_bytes);
+ uint8_t * base = data_device + (size_t) e * expert_bytes;
+
+ auto * ql_ptr = base;
+ auto * qh_ptr = ql_ptr + (QK_K / 2) * blocks_per_expert;
+ auto * scales_ptr = qh_ptr + (QK_K / 4) * blocks_per_expert;
+ auto * d_ptr = (sycl::half *) (scales_ptr + (QK_K / 16) * blocks_per_expert);
+
+ for (int j = 0; j < QK_K / 2; ++j) {
+ ql_ptr[ib * (QK_K / 2) + j] = x[ib].ql[j];
+ }
+ for (int j = 0; j < QK_K / 4; ++j) {
+ qh_ptr[ib * (QK_K / 4) + j] = x[ib].qh[j];
+ }
+ for (int j = 0; j < QK_K / 16; ++j) {
+ scales_ptr[ib * (QK_K / 16) + j] = x[ib].scales[j];
+ }
+ d_ptr[ib] = x[ib].d;
+ });
+ if (!g_ggml_sycl_use_async_mem_op) {
+ reorder_event.wait_and_throw();
+ }
+ return true;
+}
+
static bool reorder_qw_q3_k(uint8_t * data_device, size_t size, size_t offset, dpct::queue_ptr stream) {
GGML_ASSERT(size % sizeof(block_q3_K) == 0);
GGML_ASSERT(offset % sizeof(block_q3_K) == 0);
size_t nrows = src0->ne[1];
size_t size = ggml_nbytes(src0);
+ // MoE expert weights are addressed per expert via nb[2], so each slice must
+ // remain self-contained after reorder.
+ if (src0->ne[2] > 1) {
+ GGML_ASSERT((size_t) size == (size_t) src0->ne[2] * src0->nb[2]);
+ switch (src0->type) {
+ case GGML_TYPE_Q4_K:
+ return reorder_qw_q4_k_moe(data_device, src0->nb[2], src0->ne[2], stream);
+ case GGML_TYPE_Q5_K:
+ return reorder_qw_q5_k_moe(data_device, src0->nb[2], src0->ne[2], stream);
+ case GGML_TYPE_Q6_K:
+ return reorder_qw_q6_k_moe(data_device, src0->nb[2], src0->ne[2], stream);
+ default:
+ return false;
+ }
+ }
+
switch (src0->type) {
case GGML_TYPE_Q4_0:
return reorder_qw_q4_0(data_device, ncols, nrows, size, 0, stream);
case GGML_TYPE_Q6_K:
return reorder_qw_q6_k(data_device, size, 0, stream);
default:
- GGML_ABORT("reorder_qw() called with unsupported type");
return false;
}
}
}
}
+// Lazily reorder supported MoE expert weights once their fused path is used.
+static void opt_for_reorder_id(ggml_backend_sycl_context * ctx, const ggml_tensor * src0) {
+ if (g_ggml_sycl_disable_optimize || !ctx->opt_feature.reorder) {
+ return;
+ }
+ if (src0->type != GGML_TYPE_Q4_K && src0->type != GGML_TYPE_Q5_K && src0->type != GGML_TYPE_Q6_K) {
+ return;
+ }
+ ggml_tensor_extra_gpu * extra = static_cast<ggml_tensor_extra_gpu *>(src0->extra);
+ if (!extra || extra->optimized_feature.reorder) {
+ return;
+ }
+ if (reorder_qw(src0, ctx->stream())) {
+ extra->optimized_feature.reorder = true;
+ }
+}
+
static bool can_use_dequantize_mul_mat_vec(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
// The F16/BF16 qk=1 kernel iterates with stride 2*DMMV_X, requiring ne[0] to be
if (ne10 != src0->ne[0] || ne10 % QK8_1 != 0) return false;
if (!ggml_is_contiguous(src1)) return false;
- // Reorder layout not supported; fall back.
- const ggml_tensor_extra_gpu * src0_extra =
- static_cast<const ggml_tensor_extra_gpu *>(src0->extra);
- if (src0_extra && src0_extra->optimized_feature.reorder) return false;
-
const int64_t n_ids_per_group = ids->ne[0];
if (ids->ne[1] != 1) return false;
if (ne11 != 1 && ne11 != n_ids_per_group) return false;
const int n_experts_used = (int) n_ids_per_group;
const int nrows = (int) src0->ne[1];
+ // Lazily reorder the (Q4_K) expert weights into a per-expert SoA layout, then run the reorder
+ // GEMV. Placed after the bail checks so a non-dispatchable op does not pay the reorder cost.
+ opt_for_reorder_id(&ctx, src0);
+ const ggml_tensor_extra_gpu * src0_extra =
+ static_cast<const ggml_tensor_extra_gpu *>(src0->extra);
+ const bool use_reorder = src0_extra && src0_extra->optimized_feature.reorder;
+
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 *) src1->data, src1_ddq, (int) ne10, (int) ne11,
- src1_padded_cols, stream);
+ if (use_reorder) {
+ quantize_row_q8_1_sycl<quantize_and_reorder_q8_1_soa>(
+ (const float *) src1->data, src1_ddq, (int) ne10, (int) ne11,
+ src1_padded_cols, stream);
+ } else {
+ quantize_row_q8_1_sycl<quantize_q8_1>(
+ (const float *) src1->data, src1_ddq, (int) ne10, (int) ne11,
+ src1_padded_cols, stream);
+ }
const size_t bytes_per_qrow = (size_t) src1_padded_cols * sizeof(block_q8_1) / QK8_1;
const size_t src1_row_stride = (ne11 == 1) ? 0 : bytes_per_qrow;
+ if (use_reorder) {
+ return ggml_sycl_mul_mat_vec_q_id_reorder(
+ src0->type, src0->data, src1_ddq, (const int32_t *) ids->data,
+ (float *) dst->data, (int) ne10, nrows, n_experts_used,
+ /*expert_weight_stride=*/ src0->nb[2],
+ /*dst_row_stride=*/ dst->nb[1],
+ src1_row_stride, stream);
+ }
return ggml_sycl_mul_mat_vec_q_id(
src0->type, src0->data, src1_ddq, (const int32_t *) ids->data,
(float *) dst->data, (int) ne10, nrows, n_experts_used,
return false;
}
}
+
+// Reorder (SoA) MoE expert GEMV: MoE expert/row/lane indexing (from mul_mat_vec_q_moe) with the
+// dense-reorder per-block reads (from mul_mat_vec_q_reorder). Each expert slice in vx_base is a
+// self-contained SoA, so nblocks = nrows*(ncols/qk) per expert and the constant expert stride holds.
+template <typename reorder_vec_dot_q_sycl>
+static void mul_mat_vec_q_moe_reorder(
+ const void * __restrict__ vx_base, const void * __restrict__ vy_base,
+ float * __restrict__ dst_base, const int32_t * __restrict__ ids_dev,
+ const int ncols, const int nrows,
+ const size_t expert_weight_stride, const size_t dst_row_stride,
+ const size_t src1_row_stride,
+ const sycl::nd_item<3> & item_ct1) {
+ using block_type = ggml_sycl_reordered::block_q_t<reorder_vec_dot_q_sycl::gtype>;
+ using block_traits = typename block_type::traits;
+
+ const int expert_idx = item_ct1.get_group(1);
+ const int i02 = ids_dev[expert_idx];
+
+ const char * vx = (const char *) vx_base + (size_t) i02 * expert_weight_stride;
+ const char * vy = (const char *) vy_base + (size_t) expert_idx * src1_row_stride;
+ float * dst = (float *) ((char *) dst_base + (size_t) expert_idx * dst_row_stride);
+
+ 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 auto sg = item_ct1.get_sub_group();
+
+ const int blocks_per_row = ncols / block_traits::qk;
+ constexpr int blocks_per_subgroup = ceil_div(block_traits::vdr_mmvq * WARP_SIZE, block_traits::qi);
+ constexpr int block_elements_per_subgroup = block_traits::qi / block_traits::vdr_mmvq;
+ const int nblocks = nrows * (ncols / block_traits::qk);
+
+ static_assert(blocks_per_subgroup > 0);
+ static_assert(block_elements_per_subgroup > 0);
+
+ float partial_sum = 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;
+
+ 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();
+ const int8_t * q8_1_quant_ptr = (const int8_t *) vy + iby * QK8_1;
+ const sycl::half2 * q8_1_ds_ptr = (const sycl::half2 *) ((const char *) vy + ncols + iby * sizeof(sycl::half2));
+
+#pragma unroll
+ for (int elem = 0; elem < block_elements_per_subgroup; elem += WARP_SIZE) {
+ const int iqs = elem + block_traits::vdr_mmvq * (sg.get_local_linear_id() % block_elements_per_subgroup);
+ partial_sum += reorder_vec_dot_q_sycl()(vx, bx_offset, d_offset, q8_1_quant_ptr, q8_1_ds_ptr, iqs);
+ }
+ }
+
+ auto sum = sycl::reduce_over_group(sg, partial_sum, std::plus<>());
+ if (sg.leader()) {
+ dst[row] = sum;
+ }
+}
+
+template <typename reorder_vec_dot_q_sycl>
+static void launch_mul_mat_vec_q_moe_reorder(
+ const void * vx_base, const void * vy, const int32_t * ids_dev,
+ float * dst_base, const int ncols, const int nrows, const int n_experts_used,
+ const size_t expert_weight_stride, const size_t dst_row_stride,
+ const size_t src1_row_stride,
+ dpct::queue_ptr stream) {
+ const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
+ const sycl::range<3> block_nums(1, (unsigned) n_experts_used, (unsigned) 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> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
+ mul_mat_vec_q_moe_reorder<reorder_vec_dot_q_sycl>(
+ vx_base, vy, dst_base, ids_dev, ncols, nrows,
+ expert_weight_stride, dst_row_stride, src1_row_stride, item);
+ });
+ });
+}
+
+bool ggml_sycl_mul_mat_vec_q_id_reorder(
+ enum ggml_type src0_type,
+ const void * vx_base,
+ const void * vy,
+ const int32_t * ids_dev,
+ float * dst_base,
+ int ncols,
+ int nrows,
+ int n_experts_used,
+ size_t expert_weight_stride,
+ size_t dst_row_stride,
+ size_t src1_row_stride,
+ dpct::queue_ptr stream) {
+ switch (src0_type) {
+ case GGML_TYPE_Q4_K:
+ launch_mul_mat_vec_q_moe_reorder<reorder_vec_dot_q_sycl<GGML_TYPE_Q4_K>>(
+ vx_base, vy, ids_dev, dst_base, ncols, nrows, n_experts_used,
+ expert_weight_stride, dst_row_stride, src1_row_stride, stream);
+ return true;
+ case GGML_TYPE_Q5_K:
+ launch_mul_mat_vec_q_moe_reorder<reorder_vec_dot_q_sycl<GGML_TYPE_Q5_K>>(
+ vx_base, vy, ids_dev, dst_base, ncols, nrows, n_experts_used,
+ expert_weight_stride, dst_row_stride, src1_row_stride, stream);
+ return true;
+ case GGML_TYPE_Q6_K:
+ launch_mul_mat_vec_q_moe_reorder<reorder_vec_dot_q_sycl<GGML_TYPE_Q6_K>>(
+ vx_base, vy, ids_dev, dst_base, ncols, nrows, n_experts_used,
+ expert_weight_stride, dst_row_stride, src1_row_stride, stream);
+ return true;
+ default:
+ return false;
+ }
+}
size_t src1_row_stride, // 0 = shared src1, else per-expert stride in bytes
dpct::queue_ptr stream);
+// Reorder (SoA) variant of the fused MoE expert GEMV.
+// vx_base: each expert slice (stride expert_weight_stride == src0->nb[2]) is a self-contained reorder/SoA layout.
+// vy: src1 quantized with quantize_and_reorder_q8_1_soa (per-row SoA). Returns false if src0_type isn't handled.
+bool ggml_sycl_mul_mat_vec_q_id_reorder(
+ enum ggml_type src0_type,
+ const void * vx_base,
+ const void * vy,
+ const int32_t * ids_dev,
+ float * dst_base,
+ int ncols,
+ int nrows,
+ int n_experts_used,
+ size_t expert_weight_stride,
+ size_t dst_row_stride,
+ size_t src1_row_stride,
+ dpct::queue_ptr stream);
+
#endif // GGML_SYCL_MMVQ_HPP
overview / pkg / ggml / sources / llama.cpp / commitdiff