@@ -11155,25 +11155,6 @@ static void mul_mat_vec_iq3_xxs_q8_1_sycl(const void *vx, const void *vy,
11155
11155
}
11156
11156
11157
11157
11158
-
11159
- template <int qk, int qi, typename block_q_t, int vdr,
11160
- vec_dot_q_sycl_t vec_dot_q_sycl>
11161
- static void mul_mat_vec_q_sycl_submitter(const void *vx, const void *vy,
11162
- float *dst, const int ncols,
11163
- const int nrows,
11164
- dpct::queue_ptr stream) {
11165
- GGML_ASSERT(ncols % QK4_0 == 0);
11166
- const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
11167
- const sycl::range<3> block_nums(1, 1, block_num_y);
11168
- const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
11169
- stream->parallel_for(
11170
- sycl::nd_range<3>(block_nums * block_dims, block_dims), [=
11171
- ](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] {
11172
- mul_mat_vec_q<qk, qi, block_q_t, vdr, vec_dot_q_sycl>(
11173
- vx, vy, dst, ncols, nrows, item_ct1);
11174
- });
11175
- }
11176
-
11177
11158
static void ggml_mul_mat_q4_0_q8_1_sycl(const void *vx, const void *vy,
11178
11159
float *dst, const int ncols_x,
11179
11160
const int nrows_x, const int ncols_y,
0 commit comments