Skip to content

Commit 5d2bf08

Browse files
qnixsynapseqnixsynapse
committed
SYCL: Implement few same quantized type copy kernels
1 parent d74e94c commit 5d2bf08

File tree

2 files changed

+152
-0
lines changed

2 files changed

+152
-0
lines changed

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

Lines changed: 137 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3,6 +3,8 @@
33
#include <float.h>
44

55
#include "dequantize.hpp"
6+
#include "ggml-sycl/common.hpp"
7+
#include "ggml.h"
68

79
static __dpct_inline__ int best_index_int8(int n, const int8_t * val, float x) {
810
if (x <= val[0]) {
@@ -116,6 +118,38 @@ static void cpy_blck_f32_q8_0(const char * cxi, char * cdsti) {
116118
}
117119
}
118120

121+
/* quantized type same copy */
122+
static void cpy_block_q8_0_q8_0(const char * cxi, char * cdsti) {
123+
const block_q8_0 * xi = (const block_q8_0 *) cxi;
124+
block_q8_0 * dsti = (block_q8_0 *) cdsti;
125+
*dsti = *xi;
126+
}
127+
128+
static void cpy_block_q5_0_q5_0(const char * cxi, char * cdsti) {
129+
const block_q5_0 * xi = (const block_q5_0 *) cxi;
130+
block_q5_0 * dsti = (block_q5_0 *) cdsti;
131+
*dsti = *xi;
132+
}
133+
134+
135+
static void cpy_block_q5_1_q5_1(const char * cxi, char * cdsti) {
136+
const block_q5_1 * xi = (const block_q5_1 *) cxi;
137+
block_q5_1 * dsti = (block_q5_1 *) cdsti;
138+
*dsti = *xi;
139+
}
140+
141+
static void cpy_block_q4_0_q4_0(const char * cxi, char * cdsti) {
142+
const block_q4_0 * xi = (const block_q4_0 *) cxi;
143+
block_q4_0 * dsti = (block_q4_0 *) cdsti;
144+
*dsti = *xi;
145+
}
146+
147+
static void cpy_block_q4_1_q4_1(const char * cxi, char * cdsti) {
148+
const block_q4_1 * xi = (const block_q4_1 *) cxi;
149+
block_q4_1 * dsti = (block_q4_1 *) cdsti;
150+
*dsti = *xi;
151+
}
152+
119153
static void cpy_blck_q8_0_f32(const char * cxi, char * cdsti) {
120154
float * cdstf = (float *) (cdsti);
121155

@@ -311,6 +345,34 @@ template <dequantize_kernel_t dequant, int qk> static void cpy_blck_q_f32(const
311345
}
312346
}
313347

348+
349+
template <cpy_kernel_t cpy_blck, int qk>
350+
static void cpy_q_q(const char * cx, char * cdst, const int ne, const int ne00, const int ne01, const int ne02,
351+
const int nb00, const int nb01, const int nb02, const int nb03, const int ne10, const int ne11,
352+
const int ne12, const int nb10, const int nb11, const int nb12, const int nb13,
353+
const sycl::nd_item<3> & item_ct1) {
354+
const int i = (item_ct1.get_local_range(2) * item_ct1.get_group(2) + item_ct1.get_local_id(2)) * qk;
355+
356+
if (i >= ne) {
357+
return;
358+
}
359+
360+
const int i03 = i / (ne00 * ne01 * ne02);
361+
const int i02 = (i - i03 * ne00 * ne01 * ne02) / (ne00 * ne01);
362+
const int i01 = (i - i03 * ne00 * ne01 * ne02 - i02 * ne01 * ne00) / ne00;
363+
const int i00 = i - i03 * ne00 * ne01 * ne02 - i02 * ne01 * ne00 - i01 * ne00;
364+
const int x_offset = (i00 / qk) * nb00 + i01 * nb01 + i02 * nb02 + i03 * nb03;
365+
366+
367+
const int i13 = i / (ne10 * ne11 * ne12);
368+
const int i12 = (i - i13 * ne10 * ne11 * ne12) / (ne10 * ne11);
369+
const int i11 = (i - i13 * ne10 * ne11 * ne12 - i12 * ne10 * ne11) / ne10;
370+
const int i10 = i - i13 * ne10 * ne11 * ne12 - i12 * ne10 * ne11 - i11 * ne10;
371+
const int dst_offset = (i10 / qk) * nb10 + i11 * nb11 + i12 * nb12 + i13 * nb13;
372+
373+
cpy_blck(cx + x_offset, cdst + dst_offset);
374+
}
375+
314376
template <cpy_kernel_t cpy_blck, int qk>
315377
static void cpy_f32_q(const char * cx, char * cdst, const int ne, const int ne00, const int ne01, const int ne02,
316378
const int nb00, const int nb01, const int nb02, const int nb03, const int ne10, const int ne11,
@@ -322,6 +384,7 @@ static void cpy_f32_q(const char * cx, char * cdst, const int ne, const int ne00
322384
return;
323385
}
324386

387+
325388
const int i03 = i / (ne00 * ne01 * ne02);
326389
const int i02 = (i - i03 * ne00 * ne01 * ne02) / (ne00 * ne01);
327390
const int i01 = (i - i03 * ne00 * ne01 * ne02 - i02 * ne01 * ne00) / ne00;
@@ -615,6 +678,70 @@ static void ggml_cpy_i32_i32_sycl(const char * cx, char * cdst, const int ne, co
615678
}
616679
}
617680

681+
static void ggml_cpy_q8_0_q8_0(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
682+
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
683+
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
684+
const int nb12, const int nb13, queue_ptr stream) {
685+
686+
const int num_blocks = ne;
687+
stream->parallel_for(
688+
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)), [=](sycl::nd_item<3> item_ct1) {
689+
cpy_q_q<cpy_block_q8_0_q8_0, QK8_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
690+
});
691+
}
692+
693+
694+
static void ggml_cpy_q5_0_q5_0(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
695+
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
696+
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
697+
const int nb12, const int nb13, queue_ptr stream) {
698+
699+
const int num_blocks = ne;
700+
stream->parallel_for(
701+
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)), [=](sycl::nd_item<3> item_ct1) {
702+
cpy_q_q<cpy_block_q5_0_q5_0, QK5_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
703+
});
704+
}
705+
706+
707+
static void ggml_cpy_q5_1_q5_1(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
708+
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
709+
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
710+
const int nb12, const int nb13, queue_ptr stream) {
711+
712+
const int num_blocks = ne;
713+
stream->parallel_for(
714+
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)), [=](sycl::nd_item<3> item_ct1) {
715+
cpy_q_q<cpy_block_q5_1_q5_1, QK5_1>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
716+
});
717+
}
718+
719+
720+
static void ggml_cpy_q4_0_q4_0(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
721+
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
722+
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
723+
const int nb12, const int nb13, queue_ptr stream) {
724+
725+
const int num_blocks = ne;
726+
stream->parallel_for(
727+
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)), [=](sycl::nd_item<3> item_ct1) {
728+
cpy_q_q<cpy_block_q4_0_q4_0, QK4_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
729+
});
730+
}
731+
732+
733+
static void ggml_cpy_q4_1_q4_1(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
734+
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
735+
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
736+
const int nb12, const int nb13, queue_ptr stream) {
737+
738+
const int num_blocks = ne;
739+
stream->parallel_for(
740+
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)), [=](sycl::nd_item<3> item_ct1) {
741+
cpy_q_q<cpy_block_q4_1_q4_1, QK4_1>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
742+
});
743+
}
744+
618745
void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1) try {
619746
// Unlike other operators ggml_sycl_cpy takes 2 distinct tensors instead of a dst ggml_tensor and rely on its src field
620747
scope_op_debug_print scope_dbg_print(__func__, src1, /*num_src=*/0,
@@ -684,6 +811,16 @@ void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, co
684811
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_IQ4_NL) {
685812
ggml_cpy_f32_iq4_nl_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
686813
nb10, nb11, nb12, nb13, main_stream);
814+
} else if (src0->type == GGML_TYPE_Q8_0 && src1->type == GGML_TYPE_Q8_0) {
815+
ggml_cpy_q8_0_q8_0(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
816+
} else if (src0->type == GGML_TYPE_Q5_0 && src1->type == GGML_TYPE_Q5_0) {
817+
ggml_cpy_q5_0_q5_0(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
818+
} else if (src0->type == GGML_TYPE_Q5_1 && src1->type == GGML_TYPE_Q5_1) {
819+
ggml_cpy_q5_1_q5_1(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
820+
} else if (src0->type == GGML_TYPE_Q4_0 && src1->type == GGML_TYPE_Q4_0) {
821+
ggml_cpy_q4_0_q4_0(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
822+
} else if (src0->type == GGML_TYPE_Q4_1 && src1->type == GGML_TYPE_Q4_1) {
823+
ggml_cpy_q4_1_q4_1(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
687824
} else {
688825
GGML_LOG_ERROR("%s: unsupported type combination (%s to %s)\n", __func__, ggml_type_name(src0->type),
689826
ggml_type_name(src1->type));

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

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4206,6 +4206,21 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g
42064206
if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_IQ4_NL) {
42074207
return true;
42084208
}
4209+
if(src0_type == GGML_TYPE_Q8_0 && src1_type == GGML_TYPE_Q8_0) {
4210+
return true;
4211+
}
4212+
if(src0_type == GGML_TYPE_Q5_0 && src1_type == GGML_TYPE_Q5_0) {
4213+
return true;
4214+
}
4215+
if(src0_type == GGML_TYPE_Q5_1 && src1_type == GGML_TYPE_Q5_1) {
4216+
return true;
4217+
}
4218+
if(src0_type == GGML_TYPE_Q4_0 && src1_type == GGML_TYPE_Q4_0) {
4219+
return true;
4220+
}
4221+
if(src0_type == GGML_TYPE_Q4_1 && src1_type == GGML_TYPE_Q4_1) {
4222+
return true;
4223+
}
42094224
return false;
42104225
}
42114226
case GGML_OP_CONCAT:

0 commit comments

Comments
 (0)