Skip to content

Commit f99c653

Browse files
authored
sycl-exp : Revert "Minor arithmetic improvement to mmvq wrapper kernel (#7172)" (#7980)
* Revert "Minor arithmetic improvement to mmvq wrapper kernel (#7172)" This reverts commit 8c570c9. * Update ggml-sycl.cpp
1 parent e9aa742 commit f99c653

File tree

1 file changed

+8
-10
lines changed

1 file changed

+8
-10
lines changed

ggml-sycl.cpp

Lines changed: 8 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -7984,26 +7984,24 @@ static void mul_mat_vec_q(const void * __restrict__ vx, const void * __restrict_
79847984
const int blocks_per_row = ncols / qk;
79857985
const int blocks_per_warp = vdr * WARP_SIZE / qi;
79867986

7987-
const int qi_vdr = (qi / vdr); // N_threads processing 1 qk block
7988-
79897987
// partial sum for each thread
79907988
float tmp = 0.0f;
79917989

79927990
const block_q_t * x = (const block_q_t *) vx;
79937991
const block_q8_1 * y = (const block_q8_1 *) vy;
79947992

7995-
for (int i = item_ct1.get_local_id(2) / qi_vdr; i < blocks_per_row;
7993+
for (int i = item_ct1.get_local_id(2) / (qi / vdr); i < blocks_per_row;
79967994
i += blocks_per_warp) {
7997-
const int ibx = row * blocks_per_row + i; // x block index
7995+
const int ibx = row*blocks_per_row + i; // x block index
79987996

7999-
const int iby = i * (qk / QK8_1); // y block index that aligns with ibx
7997+
const int iby = i * (qk/QK8_1); // y block index that aligns with ibx
80007998

8001-
const int iqs =
8002-
vdr *
8003-
(item_ct1.get_local_id(2) -
8004-
i * qi_vdr); // x block quant index when casting the quants to int
7999+
const int iqs =
8000+
vdr *
8001+
(item_ct1.get_local_id(2) %
8002+
(qi / vdr)); // x block quant index when casting the quants to int
80058003

8006-
tmp += vec_dot_q_sycl(&x[ibx], &y[iby], iqs);
8004+
tmp += vec_dot_q_sycl(&x[ibx], &y[iby], iqs);
80078005
}
80088006

80098007
// sum up partial sums and write back result

0 commit comments

Comments
 (0)