Skip to content

Repository navigation

Commit 3bb63cb

Browse files
skipzx-labSam
skipzx-lab
and
Sam
authored
sycl : fix get_rows on reordered quantized weights (ggml-org#30430)
Co-authored-by: Sam <sam@garagellm.tail7dd858.ts.net>
1 parent f711c8f commit 3bb63cb

4 files changed

Lines changed: 262 additions & 21 deletions

File tree

‎ggml/src/ggml-sycl/convert.cpp‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -85,7 +85,7 @@ static void dequantize_row_q2_K_sycl_reorder(const void *vx, dst_t *y, const int
8585
stream->parallel_for(
8686
sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
8787
[=](sycl::nd_item<3> item_ct1) {
88-
dequantize_block_q2_K_reorder(vx, y, item_ct1, nb);
88+
dequantize_block_q2_K_reorder(vx, y + item_ct1.get_group(2) * QK_K, item_ct1.get_group(2), item_ct1, nb);
8989
});
9090
}
9191

@@ -129,7 +129,7 @@ static void dequantize_row_q3_K_sycl_reorder(const void *vx, dst_t *y, const int
129129
stream->parallel_for(
130130
sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
131131
[=](sycl::nd_item<3> item_ct1) {
132-
dequantize_block_q3_K_reorder(vx, y, item_ct1, nb);
132+
dequantize_block_q3_K_reorder(vx, y + item_ct1.get_group(2) * QK_K, item_ct1.get_group(2), item_ct1, nb);
133133
});
134134
}
135135

@@ -231,7 +231,7 @@ static void dequantize_row_q4_K_sycl_reorder(const void * vx, dst_t * y, const i
231231

232232
cgh.parallel_for(sycl::nd_range<1>(sycl::range<1>(global_size), sycl::range<1>(local_size)),
233233
[=](sycl::nd_item<1> item_ct1) {
234-
dequantize_block_q4_K_reorder(vx, y, get_pointer(scale_local_acc), item_ct1, nb);
234+
dequantize_block_q4_K_reorder(vx, y + item_ct1.get_group(0) * QK_K, item_ct1.get_group(0), get_pointer(scale_local_acc), item_ct1, nb);
235235
});
236236
});
237237
}
@@ -258,7 +258,7 @@ static void dequantize_row_q5_K_sycl_reorder(const void * vx, dst_t * y, const i
258258
cgh.parallel_for(
259259
sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
260260
[=](sycl::nd_item<3> item_ct1) {
261-
dequantize_block_q5_K_reorder(vx, y, get_pointer(scale_local_acc), item_ct1, nb);
261+
dequantize_block_q5_K_reorder(vx, y + item_ct1.get_group(2) * QK_K, item_ct1.get_group(2), get_pointer(scale_local_acc), item_ct1, nb);
262262
});
263263
});
264264
}
@@ -303,7 +303,7 @@ static void dequantize_row_q6_K_sycl_reorder(const void * vx, dst_t * y, const i
303303

304304
stream->parallel_for(
305305
sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
306-
[=](sycl::nd_item<3> item_ct1) { dequantize_block_q6_K_reorder(vx, y, item_ct1, nb); });
306+
[=](sycl::nd_item<3> item_ct1) { dequantize_block_q6_K_reorder(vx, y + item_ct1.get_group(2) * QK_K, item_ct1.get_group(2), item_ct1, nb); });
307307
}
308308

309309
template <typename dst_t>

‎ggml/src/ggml-sycl/dequantize.hpp‎

Lines changed: 13 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -926,10 +926,9 @@ static void dequantize_block_q2_K(const void * __restrict__ vx, dst_t * __restri
926926
}
927927

928928
template<typename dst_t>
929-
static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy,
929+
static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t i,
930930
const sycl::nd_item<3> & item_ct1, int64_t n_blocks) {
931931
#if QK_K == 256
932-
const int64_t i = item_ct1.get_group(2);
933932
if (i >= n_blocks) {
934933
return;
935934
}
@@ -949,7 +948,7 @@ static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t *
949948
const int64_t is = 8 * n + l / 16;
950949

951950
const uint8_t q = qs[32 * n + l];
952-
dst_t * y = yy + i * QK_K + 128 * n;
951+
dst_t * y = yy + 128 * n;
953952

954953
const float dall = (*dm)[0];
955954
const float dmin = (*dm)[1];
@@ -960,6 +959,7 @@ static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t *
960959
#else
961960
GGML_UNUSED(vx);
962961
GGML_UNUSED(yy);
962+
GGML_UNUSED(i);
963963
GGML_UNUSED(item_ct1);
964964
GGML_UNUSED(n_blocks);
965965
GGML_ABORT("Q2_K reorder dequantize not supported for QK_K != 256");
@@ -1022,10 +1022,9 @@ static void dequantize_block_q3_K(const void * __restrict__ vx, dst_t * __restri
10221022
}
10231023

10241024
template<typename dst_t>
1025-
static void dequantize_block_q3_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy,
1025+
static void dequantize_block_q3_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t i,
10261026
const sycl::nd_item<3> & item_ct1, int64_t n_blocks) {
10271027
#if QK_K == 256
1028-
const int64_t i = item_ct1.get_group(2);
10291028
if (i >= n_blocks) {
10301029
return;
10311030
}
@@ -1062,7 +1061,7 @@ static void dequantize_block_q3_K_reorder(const void * __restrict__ vx, dst_t *
10621061

10631062
const float dl = d_all * (us - 32);
10641063

1065-
dst_t * y = yy + i * QK_K + 128 * n + 32 * j;
1064+
dst_t * y = yy + 128 * n + 32 * j;
10661065
const uint8_t * q = qs + 32 * n;
10671066
const uint8_t * hm = hmask;
10681067

@@ -1072,6 +1071,7 @@ static void dequantize_block_q3_K_reorder(const void * __restrict__ vx, dst_t *
10721071
#else
10731072
GGML_UNUSED(vx);
10741073
GGML_UNUSED(yy);
1074+
GGML_UNUSED(i);
10751075
GGML_UNUSED(item_ct1);
10761076
GGML_UNUSED(n_blocks);
10771077
GGML_ABORT("Q3_K reorder dequantize not supported for QK_K != 256");
@@ -1160,14 +1160,13 @@ inline void dequantize_q4_K_common(dst_t * __restrict__ y, const uint8_t * __res
11601160
}
11611161

11621162
template <typename dst_t>
1163-
static void dequantize_block_q4_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, uint8_t * scales_local,
1164-
const sycl::nd_item<1> & item_ct1, int64_t nb) {
1165-
const int64_t i = item_ct1.get_group(0); // block index
1163+
static void dequantize_block_q4_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t i,
1164+
uint8_t * scales_local, const sycl::nd_item<1> & item_ct1, int64_t nb) {
11661165
const int64_t tid = item_ct1.get_local_id(0); // thread index within block
11671166
const int64_t il = tid / 8;
11681167
const int64_t ir = tid % 8;
11691168

1170-
dst_t * y = yy + i * QK_K + 64 * il + 4 * ir;
1169+
dst_t * y = yy + 64 * il + 4 * ir;
11711170

11721171
const uint8_t * base = static_cast<const uint8_t *>(vx);
11731172
const size_t qs_offset = i * (QK_K / 2);
@@ -1237,9 +1236,8 @@ static void dequantize_block_q5_K_wide(const void * __restrict__ vx, dst_t * __r
12371236
}
12381237

12391238
template <typename dst_t>
1240-
static void dequantize_block_q5_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy,
1239+
static void dequantize_block_q5_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t ib,
12411240
uint8_t * scales_local, const sycl::nd_item<3> & item_ct1, int64_t n_blocks) {
1242-
const int64_t ib = item_ct1.get_group(2);
12431241

12441242
#if QK_K == 256
12451243
// assume 64 threads
@@ -1248,7 +1246,7 @@ static void dequantize_block_q5_K_reorder(const void * __restrict__ vx, dst_t *
12481246
const int64_t ir = tid % 16; // 0...15
12491247
const int64_t is = 2 * il;
12501248

1251-
dst_t * y = yy + ib * QK_K + 64 * il + 2 * ir;
1249+
dst_t * y = yy + 64 * il + 2 * ir;
12521250

12531251
const uint8_t * base = static_cast<const uint8_t *>(vx);
12541252

@@ -1340,9 +1338,8 @@ static void dequantize_block_q6_K(const void * __restrict__ vx, dst_t * __restri
13401338
}
13411339

13421340
template <typename dst_t>
1343-
static void dequantize_block_q6_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy,
1341+
static void dequantize_block_q6_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t ib,
13441342
const sycl::nd_item<3> & item_ct1, int64_t n_blocks) {
1345-
const int64_t ib = item_ct1.get_group(2);
13461343

13471344
const int64_t tid = item_ct1.get_local_id(2);
13481345
const int64_t ip = tid / 32; // ip is 0 or 1
@@ -1359,7 +1356,7 @@ static void dequantize_block_q6_K_reorder(const void * __restrict__ vx, dst_t *
13591356
const uint8_t * scales_ptr = base_ptr + base_scales_offset;
13601357
const ggml_half * d = (const ggml_half *) (base_ptr + base_d_offset) + ib;
13611358

1362-
dst_t * y = yy + ib * QK_K + 128 * ip + il;
1359+
dst_t * y = yy + 128 * ip + il;
13631360

13641361
const uint8_t * ql = ql_ptr + 64 * ip + il;
13651362
const uint8_t qh = *(qh_ptr + 32 * ip + il);

0 commit comments

Comments
 (0)