Skip to content

Commit 3ec8eb3

Browse files
s-Nicks-Nick
committed
Add Reorder to Q6_K mmvq implementation
Signed-off-by: nscipione <[email protected]>
1 parent 26b79b6 commit 3ec8eb3

File tree

6 files changed

+242
-25
lines changed

6 files changed

+242
-25
lines changed

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

Lines changed: 21 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -265,6 +265,17 @@ static void dequantize_row_q6_K_sycl(const void *vx, dst_t *y, const int64_t k,
265265
#endif
266266
}
267267

268+
template <typename dst_t>
269+
static void dequantize_row_q6_K_sycl_reorder(const void * vx, dst_t * y, const int64_t k, dpct::queue_ptr stream) {
270+
const int64_t nb = k / QK_K;
271+
272+
dpct::has_capability_or_fail(stream->get_device(), { sycl::aspect::fp16 });
273+
274+
stream->parallel_for(
275+
sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
276+
[=](sycl::nd_item<3> item_ct1) { dequantize_block_q6_K_reorder(vx, y, item_ct1, nb); });
277+
}
278+
268279
template <typename dst_t>
269280
static void dequantize_row_iq1_s_sycl(const void *vx, dst_t *y, const int64_t k,
270281
dpct::queue_ptr stream) {
@@ -530,7 +541,11 @@ to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst) {
530541
case GGML_TYPE_Q5_K:
531542
return dequantize_row_q5_K_sycl;
532543
case GGML_TYPE_Q6_K:
533-
return dequantize_row_q6_K_sycl;
544+
if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
545+
return dequantize_row_q6_K_sycl_reorder;
546+
} else {
547+
return dequantize_row_q6_K_sycl;
548+
}
534549
case GGML_TYPE_IQ1_S:
535550
return dequantize_row_iq1_s_sycl;
536551
case GGML_TYPE_IQ1_M:
@@ -587,7 +602,11 @@ to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor *dst) {
587602
case GGML_TYPE_Q5_K:
588603
return dequantize_row_q5_K_sycl;
589604
case GGML_TYPE_Q6_K:
590-
return dequantize_row_q6_K_sycl;
605+
if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
606+
return dequantize_row_q6_K_sycl_reorder;
607+
} else {
608+
return dequantize_row_q6_K_sycl;
609+
}
591610
case GGML_TYPE_IQ1_S:
592611
return dequantize_row_iq1_s_sycl;
593612
case GGML_TYPE_IQ1_M:

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

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -538,6 +538,38 @@ static void dequantize_block_q6_K(const void * __restrict__ vx, dst_t * __restri
538538
#endif
539539
}
540540

541+
template <typename dst_t>
542+
static void dequantize_block_q6_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy,
543+
const sycl::nd_item<3> & item_ct1, int64_t n_blocks) {
544+
const int64_t ib = item_ct1.get_group(2);
545+
546+
const int64_t tid = item_ct1.get_local_id(2);
547+
const int64_t ip = tid / 32; // ip is 0 or 1
548+
const int64_t il = tid - 32 * ip; // 0...32
549+
const int64_t is = 8 * ip + il / 16;
550+
551+
const uint8_t * base_ptr = static_cast<const uint8_t *>(vx);
552+
const auto ql_offset = ib * (QK_K / 2);
553+
const auto qh_offset = (QK_K / 2) * n_blocks + (QK_K / 4) * ib;
554+
const auto base_scales_offset = (QK_K / 2) * n_blocks + (QK_K / 4) * n_blocks + (QK_K / 16) * ib;
555+
const auto base_d_offset = ((QK_K / 2) + (QK_K / 4) + (QK_K / 16)) * n_blocks;
556+
const uint8_t * ql_ptr = base_ptr + ql_offset;
557+
const uint8_t * qh_ptr = base_ptr + qh_offset;
558+
const uint8_t * scales_ptr = base_ptr + base_scales_offset;
559+
const ggml_half * d = (const ggml_half *) (base_ptr + base_d_offset) + ib;
560+
561+
dst_t * y = yy + ib * QK_K + 128 * ip + il;
562+
563+
const uint8_t * ql = ql_ptr + 64 * ip + il;
564+
const uint8_t qh = *(qh_ptr + 32 * ip + il);
565+
const int8_t * sc = reinterpret_cast<const int8_t *>(scales_ptr + is);
566+
567+
y[0] = *d * sc[0] * ((int8_t) ((ql[0] & 0xF) | (((qh >> 0) & 3) << 4)) - 32);
568+
y[32] = *d * sc[2] * ((int8_t) ((ql[32] & 0xF) | (((qh >> 2) & 3) << 4)) - 32);
569+
y[64] = *d * sc[4] * ((int8_t) ((ql[0] >> 4) | (((qh >> 4) & 3) << 4)) - 32);
570+
y[96] = *d * sc[6] * ((int8_t) ((ql[32] >> 4) | (((qh >> 6) & 3) << 4)) - 32);
571+
}
572+
541573
template<typename dst_t>
542574
static void dequantize_block_iq2_xxs(const void * __restrict__ vx, dst_t * __restrict__ yy,
543575
const sycl::nd_item<3> &item_ct1,

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

Lines changed: 52 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -354,7 +354,8 @@ ggml_backend_sycl_buffer_init_tensor(ggml_backend_buffer_t buffer,
354354
assert(tensor->view_src->buffer->buft == buffer->buft);
355355
return GGML_STATUS_SUCCESS;
356356
}
357-
if ((tensor->type == GGML_TYPE_Q4_0 || tensor->type == GGML_TYPE_Q4_K) && !g_ggml_sycl_disable_optimize) {
357+
if ((tensor->type == GGML_TYPE_Q4_0 || tensor->type == GGML_TYPE_Q4_K || tensor->type == GGML_TYPE_Q6_K) &&
358+
!g_ggml_sycl_disable_optimize) {
358359
ggml_tensor_extra_gpu * extra = new ggml_tensor_extra_gpu{};
359360
tensor->extra = extra;
360361
ctx->tensor_extras.push_back(extra); //used to release it when destroy ctx.
@@ -2928,6 +2929,7 @@ inline bool ggml_sycl_supports_reorder_mul_mat_sycl(enum ggml_type type) {
29282929
case GGML_TYPE_Q4_0:
29292930
return true;
29302931
case GGML_TYPE_Q4_K:
2932+
case GGML_TYPE_Q6_K:
29312933
return !g_ggml_sycl_prioritize_dmmv;
29322934
default:
29332935
return false;
@@ -2947,6 +2949,7 @@ inline bool ggml_sycl_supports_reorder_mmvq(enum ggml_type type) {
29472949
switch (type) {
29482950
case GGML_TYPE_Q4_0:
29492951
case GGML_TYPE_Q4_K:
2952+
case GGML_TYPE_Q6_K:
29502953
return true;
29512954
default:
29522955
return false;
@@ -3031,6 +3034,51 @@ static void reorder_qw_q4_k(uint8_t * data_device, size_t size, size_t offset, d
30313034
sycl::free(tmp_buf, *stream);
30323035
}
30333036

3037+
static void reorder_qw_q6_k(uint8_t * data_device, size_t size, size_t offset, dpct::queue_ptr stream) {
3038+
GGML_ASSERT(size % sizeof(block_q6_K) == 0);
3039+
GGML_ASSERT(offset % sizeof(block_q6_K) == 0);
3040+
3041+
const int nblocks = size / sizeof(block_q6_K);
3042+
3043+
auto * tmp_buf = sycl::malloc_shared<uint8_t>(size, *stream);
3044+
SYCL_CHECK(CHECK_TRY_ERROR((*stream).memcpy(tmp_buf, data_device, size).wait()));
3045+
3046+
auto * ql_ptr = data_device;
3047+
auto * qh_ptr = ql_ptr + (QK_K / 2) * nblocks;
3048+
// scales are after all quants' bits so adding both to get correct offset
3049+
auto * scales_ptr = qh_ptr + (QK_K / 4) * nblocks;
3050+
sycl::half * dm_ptr = (sycl::half *) (scales_ptr + (QK_K / 16) * nblocks);
3051+
3052+
stream
3053+
->parallel_for(nblocks,
3054+
[=](auto i) {
3055+
const block_q6_K * x = (const block_q6_K *) tmp_buf;
3056+
const int ib = i;
3057+
3058+
const uint8_t * ql = x[ib].ql;
3059+
const uint8_t * qh = x[ib].qh;
3060+
uint8_t * base_ql_ptr = ql_ptr + (QK_K / 2) * ib;
3061+
uint8_t * base_qh_ptr = qh_ptr + (QK_K / 4) * ib;
3062+
uint8_t * base_scales_ptr = scales_ptr + (QK_K / 16) * ib;
3063+
3064+
for (int j = 0; j < QK_K / 2; ++j) {
3065+
base_ql_ptr[j] = ql[j];
3066+
}
3067+
for (int j = 0; j < QK_K / 4; ++j) {
3068+
base_qh_ptr[j] = qh[j];
3069+
}
3070+
3071+
for (int j = 0; j < QK_K / 16; ++j) {
3072+
base_scales_ptr[j] = x[ib].scales[j];
3073+
}
3074+
3075+
dm_ptr[ib] = x[ib].d;
3076+
})
3077+
.wait_and_throw();
3078+
3079+
sycl::free(tmp_buf, *stream);
3080+
}
3081+
30343082
static void reorder_qw(const ggml_tensor * src0, dpct::queue_ptr stream) {
30353083
uint8_t * data_device = (uint8_t *) src0->data;
30363084
size_t ncols = src0->ne[0];
@@ -3044,6 +3092,9 @@ static void reorder_qw(const ggml_tensor * src0, dpct::queue_ptr stream) {
30443092
case GGML_TYPE_Q4_K:
30453093
reorder_qw_q4_k(data_device, size, 0, stream);
30463094
break;
3095+
case GGML_TYPE_Q6_K:
3096+
reorder_qw_q6_k(data_device, size, 0, stream);
3097+
break;
30473098
default:
30483099
GGML_ABORT("reorder_qw() called with unsupported type");
30493100
break;

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

Lines changed: 28 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -35,9 +35,8 @@ static void mul_mat_vec_q_reorder(const void * __restrict__ vx, const void * __r
3535
for (int i = sg.get_local_linear_id() / block_elements_per_subgroup; i < blocks_per_row; i += blocks_per_subgroup) {
3636
const int ibx = row * blocks_per_row + i; // x block index
3737
// TODO: Generalize offsets, right now only works for quantizations that don't split high and low bits
38-
const int bx_offset = block_type::get_block_offset(ibx);
39-
const int d_offset = block_type::get_d_offset(nrows, ncols, ibx);
40-
38+
const auto bx_offset = block_type::get_block_offset(ibx, nblocks);
39+
const auto d_offset = block_type::get_d_offset(nrows, ncols, ibx);
4140
// Y block index that aligns with ibx
4241
const int iby = i * block_type::block_to_q8_1_ratio();
4342

@@ -785,6 +784,24 @@ static void mul_mat_vec_q5_K_q8_1_sycl(const void *vx, const void *vy,
785784
}
786785
}
787786

787+
static void reorder_mul_mat_vec_q6_k_q8_1_sycl(const void * vx, const void * vy, float * dst, const int ncols,
788+
const int nrows, dpct::queue_ptr stream) {
789+
GGML_ASSERT(ncols % QK_K == 0);
790+
const int block_num_y = ceil_div(nrows, GGML_SYCL_MMV_Y);
791+
constexpr size_t num_subgroups = 16;
792+
GGML_ASSERT(block_num_y % num_subgroups == 0);
793+
794+
const sycl::range<3> global_size(1, GGML_SYCL_MMV_Y, block_num_y * WARP_SIZE);
795+
const sycl::range<3> workgroup_size(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE);
796+
797+
stream->submit([&](sycl::handler & cgh) {
798+
cgh.parallel_for(sycl::nd_range<3>(global_size, workgroup_size),
799+
[=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
800+
mul_mat_vec_q_reorder<reorder_vec_dot_q_sycl<GGML_TYPE_Q6_K>>(vx, vy, dst, ncols, nrows,
801+
nd_item);
802+
});
803+
});
804+
}
788805
static void mul_mat_vec_q6_K_q8_1_sycl(const void *vx, const void *vy,
789806
float *dst, const int ncols,
790807
const int nrows,
@@ -1070,7 +1087,14 @@ void ggml_sycl_op_mul_mat_vec_q(ggml_backend_sycl_context & ctx, const ggml_tens
10701087
mul_mat_vec_q5_K_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream);
10711088
break;
10721089
case GGML_TYPE_Q6_K:
1073-
mul_mat_vec_q6_K_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream);
1090+
if ((ggml_tensor_extra_gpu *) dst->src[0]->extra &&
1091+
((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
1092+
GGML_SYCL_DEBUG("Calling reorder_mul_mat_vec_q6_k_q8_1_sycl\n");
1093+
reorder_mul_mat_vec_q6_k_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream);
1094+
} else {
1095+
GGML_SYCL_DEBUG("Calling mul_mat_vec_q6_k_q8_1_sycl\n");
1096+
mul_mat_vec_q6_K_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream);
1097+
}
10741098
break;
10751099
case GGML_TYPE_IQ1_S:
10761100
mul_mat_vec_iq1_s_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream);

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

Lines changed: 38 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -14,12 +14,13 @@
1414
#ifndef GGML_SYCL_QUANTS_HPP
1515
#define GGML_SYCL_QUANTS_HPP
1616

17+
#include <utility>
18+
1719
#include "ggml-common.h"
1820
#include "ggml.h"
1921

2022
namespace ggml_sycl_reordered {
2123

22-
2324
// The reordered block moves quants (qs) and scales(d) to two
2425
// uniform regions of memory that is contiguous in the same tensor.
2526
// What this means is that instead of having:
@@ -32,7 +33,6 @@ namespace ggml_sycl_reordered {
3233

3334
template <ggml_type type> struct block_q_t;
3435

35-
3636
// qk number of weights / quants in a block
3737
// qr number of weights in a byte (described as 'before dequantization')
3838
// for quantization types that has low and high bits split, qr is calculated with
@@ -47,10 +47,12 @@ template <> struct block_q_t<GGML_TYPE_Q4_0> {
4747
static constexpr uint32_t vdr_mmvq = 2;
4848
};
4949

50-
static constexpr int get_block_offset(const int block_index) { return block_index * (traits::qk / traits::qr); }
50+
static constexpr std::pair<int, int> get_block_offset(const int block_index, const int /* nblocks */) {
51+
return { block_index * (traits::qk / traits::qr), 0 };
52+
}
5153

52-
static constexpr int get_d_offset(int nrows, int ncols, const int block_index) {
53-
return (ncols / traits::qr * nrows) + block_index * sizeof(ggml_half);
54+
static constexpr std::pair<int, int> get_d_offset(int nrows, int ncols, const int block_index) {
55+
return { (ncols / traits::qr * nrows) + block_index * sizeof(ggml_half), 0 };
5456
}
5557

5658
static constexpr int block_to_q8_1_ratio() { return traits::qk / QK8_1; }
@@ -64,20 +66,47 @@ template <> struct block_q_t<GGML_TYPE_Q4_K> {
6466
static constexpr uint32_t vdr_mmvq = 2;
6567
};
6668

67-
static constexpr int get_block_offset(const int block_index) { return block_index * (traits::qk / traits::qr); }
69+
static constexpr std::pair<int, int> get_block_offset(const int block_index, const int /* nblocks */) {
70+
return { block_index * (traits::qk / traits::qr), 0 };
71+
}
6872

69-
static constexpr int get_d_offset(int nrows, int ncols, const int block_index) {
73+
static constexpr std::pair<int, int> get_d_offset(int nrows, int ncols, const int block_index) {
7074
auto nblocks = (nrows * (ncols / traits::qk));
71-
return (nblocks * QK_K / 2) + (nblocks * K_SCALE_SIZE) + (block_index * sizeof(ggml_half2));
75+
return { (nblocks * QK_K / 2) + (nblocks * K_SCALE_SIZE) + (block_index * sizeof(ggml_half2)), 0 };
7276
}
7377

7478
static constexpr int block_to_q8_1_ratio() { return traits::qk / QK8_1; }
7579

7680
constexpr size_t get_total_qs_bytes(int nblocks) { return nblocks * QK_K / 2; }
7781

78-
constexpr size_t get_dm_offset(int nblocks) { return get_total_qs_bytes(nblocks) + nblocks * K_SCALE_SIZE; }
82+
constexpr int get_dm_offset(int nblocks) { return get_total_qs_bytes(nblocks) + nblocks * K_SCALE_SIZE; }
7983
};
8084

85+
template <> struct block_q_t<GGML_TYPE_Q6_K> {
86+
struct traits {
87+
static constexpr uint32_t qk = QK_K;
88+
static constexpr uint32_t qi = QI6_K;
89+
static constexpr uint32_t qr = QR6_K;
90+
static constexpr uint32_t vdr_mmvq = 1;
91+
};
92+
93+
static constexpr std::pair<int, int> get_block_offset(const int block_index, const int n_blocks) {
94+
auto low_bits_index = block_index * (traits::qk / traits::qr);
95+
// the index of high bits it's after all low bits
96+
auto high_bits_index = n_blocks * (QK_K / 2) + (block_index * (QK_K / 4));
97+
return { low_bits_index, high_bits_index };
98+
}
99+
100+
static constexpr std::pair<int, int> get_d_offset(int nrows, int ncols, const int block_index) {
101+
auto nblocks = (nrows * (ncols / traits::qk));
102+
auto total_qs_bytes = nblocks * (QK_K / 2) + nblocks * (QK_K / 4);
103+
auto block_scales = total_qs_bytes + block_index * (QK_K / 16);
104+
auto sb_scale = total_qs_bytes + nblocks * (QK_K / 16);
105+
return { block_scales, sb_scale };
106+
}
107+
108+
static constexpr int block_to_q8_1_ratio() { return traits::qk / QK8_1; }
109+
};
81110
} // namespace ggml_sycl_reordered
82111

83112
#endif // GGML_SYCL_QUANTS_HPP

0 commit comments

Comments
 (0)