diff --git a/ggml/src/ggml-sycl/convert.cpp b/ggml/src/ggml-sycl/convert.cpp index 8ba26d3a9..45c60e8e5 100644 --- a/ggml/src/ggml-sycl/convert.cpp +++ b/ggml/src/ggml-sycl/convert.cpp @@ -85,7 +85,7 @@ static void dequantize_row_q2_K_sycl_reorder(const void *vx, dst_t *y, const int stream->parallel_for( sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)), [=](sycl::nd_item<3> item_ct1) { - dequantize_block_q2_K_reorder(vx, y, item_ct1, nb); + dequantize_block_q2_K_reorder(vx, y + item_ct1.get_group(2) * QK_K, item_ct1.get_group(2), item_ct1, nb); }); } @@ -129,7 +129,7 @@ static void dequantize_row_q3_K_sycl_reorder(const void *vx, dst_t *y, const int stream->parallel_for( sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)), [=](sycl::nd_item<3> item_ct1) { - dequantize_block_q3_K_reorder(vx, y, item_ct1, nb); + dequantize_block_q3_K_reorder(vx, y + item_ct1.get_group(2) * QK_K, item_ct1.get_group(2), item_ct1, nb); }); } @@ -231,7 +231,7 @@ static void dequantize_row_q4_K_sycl_reorder(const void * vx, dst_t * y, const i cgh.parallel_for(sycl::nd_range<1>(sycl::range<1>(global_size), sycl::range<1>(local_size)), [=](sycl::nd_item<1> item_ct1) { - dequantize_block_q4_K_reorder(vx, y, get_pointer(scale_local_acc), item_ct1, nb); + 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); }); }); } @@ -258,7 +258,7 @@ static void dequantize_row_q5_K_sycl_reorder(const void * vx, dst_t * y, const i cgh.parallel_for( sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)), [=](sycl::nd_item<3> item_ct1) { - dequantize_block_q5_K_reorder(vx, y, get_pointer(scale_local_acc), item_ct1, nb); + 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); }); }); } @@ -303,7 +303,7 @@ static void dequantize_row_q6_K_sycl_reorder(const void * vx, dst_t * y, const i stream->parallel_for( sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)), - [=](sycl::nd_item<3> item_ct1) { dequantize_block_q6_K_reorder(vx, y, item_ct1, nb); }); + [=](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); }); } template diff --git a/ggml/src/ggml-sycl/dequantize.hpp b/ggml/src/ggml-sycl/dequantize.hpp index 48ee651ad..f00f97ad6 100644 --- a/ggml/src/ggml-sycl/dequantize.hpp +++ b/ggml/src/ggml-sycl/dequantize.hpp @@ -926,10 +926,9 @@ static void dequantize_block_q2_K(const void * __restrict__ vx, dst_t * __restri } template -static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, +static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t i, const sycl::nd_item<3> & item_ct1, int64_t n_blocks) { #if QK_K == 256 - const int64_t i = item_ct1.get_group(2); if (i >= n_blocks) { return; } @@ -949,7 +948,7 @@ static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t * const int64_t is = 8 * n + l / 16; const uint8_t q = qs[32 * n + l]; - dst_t * y = yy + i * QK_K + 128 * n; + dst_t * y = yy + 128 * n; const float dall = (*dm)[0]; const float dmin = (*dm)[1]; @@ -960,6 +959,7 @@ static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t * #else GGML_UNUSED(vx); GGML_UNUSED(yy); + GGML_UNUSED(i); GGML_UNUSED(item_ct1); GGML_UNUSED(n_blocks); 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 } template -static void dequantize_block_q3_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, +static void dequantize_block_q3_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t i, const sycl::nd_item<3> & item_ct1, int64_t n_blocks) { #if QK_K == 256 - const int64_t i = item_ct1.get_group(2); if (i >= n_blocks) { return; } @@ -1062,7 +1061,7 @@ static void dequantize_block_q3_K_reorder(const void * __restrict__ vx, dst_t * const float dl = d_all * (us - 32); - dst_t * y = yy + i * QK_K + 128 * n + 32 * j; + dst_t * y = yy + 128 * n + 32 * j; const uint8_t * q = qs + 32 * n; const uint8_t * hm = hmask; @@ -1072,6 +1071,7 @@ static void dequantize_block_q3_K_reorder(const void * __restrict__ vx, dst_t * #else GGML_UNUSED(vx); GGML_UNUSED(yy); + GGML_UNUSED(i); GGML_UNUSED(item_ct1); GGML_UNUSED(n_blocks); 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 } template -static void dequantize_block_q4_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, uint8_t * scales_local, - const sycl::nd_item<1> & item_ct1, int64_t nb) { - const int64_t i = item_ct1.get_group(0); // block index +static void dequantize_block_q4_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t i, + uint8_t * scales_local, const sycl::nd_item<1> & item_ct1, int64_t nb) { const int64_t tid = item_ct1.get_local_id(0); // thread index within block const int64_t il = tid / 8; const int64_t ir = tid % 8; - dst_t * y = yy + i * QK_K + 64 * il + 4 * ir; + dst_t * y = yy + 64 * il + 4 * ir; const uint8_t * base = static_cast(vx); 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 } template -static void dequantize_block_q5_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, +static void dequantize_block_q5_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t ib, uint8_t * scales_local, const sycl::nd_item<3> & item_ct1, int64_t n_blocks) { - const int64_t ib = item_ct1.get_group(2); #if QK_K == 256 // assume 64 threads @@ -1248,7 +1246,7 @@ static void dequantize_block_q5_K_reorder(const void * __restrict__ vx, dst_t * const int64_t ir = tid % 16; // 0...15 const int64_t is = 2 * il; - dst_t * y = yy + ib * QK_K + 64 * il + 2 * ir; + dst_t * y = yy + 64 * il + 2 * ir; const uint8_t * base = static_cast(vx); @@ -1340,9 +1338,8 @@ static void dequantize_block_q6_K(const void * __restrict__ vx, dst_t * __restri } template -static void dequantize_block_q6_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, +static void dequantize_block_q6_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t ib, const sycl::nd_item<3> & item_ct1, int64_t n_blocks) { - const int64_t ib = item_ct1.get_group(2); const int64_t tid = item_ct1.get_local_id(2); 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 * const uint8_t * scales_ptr = base_ptr + base_scales_offset; const ggml_half * d = (const ggml_half *) (base_ptr + base_d_offset) + ib; - dst_t * y = yy + ib * QK_K + 128 * ip + il; + dst_t * y = yy + 128 * ip + il; const uint8_t * ql = ql_ptr + 64 * ip + il; const uint8_t qh = *(qh_ptr + 32 * ip + il); diff --git a/ggml/src/ggml-sycl/getrows.cpp b/ggml/src/ggml-sycl/getrows.cpp index 4e84bd222..9c17daf86 100644 --- a/ggml/src/ggml-sycl/getrows.cpp +++ b/ggml/src/ggml-sycl/getrows.cpp @@ -137,6 +137,47 @@ static void k_get_rows_float( dst_row[i00] = src0_row[i00]; } +// src0 is in reorder layout: all qs of the tensor, then all d +template +static void k_get_rows_reorder( + const void * src0_qs, const void * src0_d, const int32_t * src1, dst_t * dst, + int64_t ne00, int64_t ne12, + size_t s1, size_t s2, size_t s3, + size_t s10, size_t s11, size_t s12, + const sycl::nd_item<3> &item_ct1) { + + const int i00 = (item_ct1.get_group(2) * item_ct1.get_local_range(2) + + item_ct1.get_local_id(2)) * + 2; + const int i10 = item_ct1.get_local_range(1) * item_ct1.get_group(1) + + item_ct1.get_local_id(1); + const int i11 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) + + item_ct1.get_local_id(0)) / + ne12; + const int i12 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) + + item_ct1.get_local_id(0)) % + ne12; + + if (i00 >= ne00) { + return; + } + + const int i01 = src1[i10*s10 + i11*s11 + i12*s12]; + + dst_t * dst_row = dst + i10*s1 + i11*s2 + i12*s3; + + const int64_t ib = (int64_t) i01*(ne00/qk) + i00/qk; // block index in the whole tensor + const int iqs = (i00%qk)/qr; + const int iybs = i00 - i00%qk; + const int y_offset = qr == 1 ? 1 : qk/2; + + dfloat2 v; + dequantize_kernel(src0_d, ib, src0_qs, ib*(qk/qr) + iqs, v); + + dst_row[iybs + iqs + 0] = v.x(); + dst_row[iybs + iqs + y_offset] = v.y(); +} + template static void get_rows_sycl(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1, ggml_tensor *dst, const void *src0_dd, @@ -173,6 +214,93 @@ static void get_rows_sycl(ggml_backend_sycl_context & ctx, const ggml_tensor *sr GGML_UNUSED(ctx); } +template +static void get_rows_sycl_reorder(const ggml_tensor *src0, const ggml_tensor *src1, + ggml_tensor *dst, const void *src0_dd, + const int32_t *src1_dd, float *dst_dd, + queue_ptr stream) { + + GGML_TENSOR_BINARY_OP_LOCALS + + // reorder is done only on 2D weights + GGML_ASSERT(ne02 == 1 && ne03 == 1); + GGML_ASSERT(ne00 % 2 == 0); + + const sycl::range<3> block_dims(1, 1, SYCL_GET_ROWS_BLOCK_SIZE); + const int block_num_x = (ne00 + 2*SYCL_GET_ROWS_BLOCK_SIZE - 1) / (2*SYCL_GET_ROWS_BLOCK_SIZE); + const sycl::range<3> block_nums(ne11 * ne12, ne10, block_num_x); + + const size_t s1 = nb1 / ggml_element_size(dst); + const size_t s2 = nb2 / ggml_element_size(dst); + const size_t s3 = nb3 / ggml_element_size(dst); + + const size_t s10 = nb10 / ggml_element_size(src1); + const size_t s11 = nb11 / ggml_element_size(src1); + const size_t s12 = nb12 / ggml_element_size(src1); + + const void * src0_d = (const char *) src0_dd + ne00*ne01/qr; + + stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), + [=](sycl::nd_item<3> item_ct1) { + k_get_rows_reorder( + src0_dd, src0_d, src1_dd, dst_dd, ne00, ne12, s1, s2, + s3, s10, s11, s12, item_ct1); + }); + + GGML_UNUSED(src0); +} + +// K-quants in reorder layout: one work-group dequantizes one QK_K block of one row +template +static void get_rows_sycl_reorder_k(const ggml_tensor *src0, const ggml_tensor *src1, + ggml_tensor *dst, const int32_t *src1_dd, float *dst_dd, + queue_ptr stream, F dequantize_block) { + + GGML_TENSOR_BINARY_OP_LOCALS + + // reorder is done only on 2D weights + GGML_ASSERT(ne02 == 1 && ne03 == 1); + GGML_ASSERT(ne00 % QK_K == 0); + + const int64_t nblk = ne00 / QK_K; + const int64_t nb = nblk * ne01; // blocks in the whole tensor + const int64_t ngroups = ne10 * ne11 * ne12 * nblk; + + const size_t s1 = nb1 / ggml_element_size(dst); + const size_t s2 = nb2 / ggml_element_size(dst); + const size_t s3 = nb3 / ggml_element_size(dst); + + const size_t s10 = nb10 / ggml_element_size(src1); + const size_t s11 = nb11 / ggml_element_size(src1); + const size_t s12 = nb12 / ggml_element_size(src1); + + stream->submit([&](sycl::handler & cgh) { + sycl::local_accessor scales_local(sycl::range<1>(nlocal), cgh); + + auto kernel = [=](sycl::nd_item item_ct1) { + const int64_t g = item_ct1.get_group(dims - 1); + const int64_t b = g % nblk; + const int64_t r = g / nblk; + const int64_t i10 = r % ne10; + const int64_t i11 = (r / ne10) % ne11; + const int64_t i12 = r / (ne10 * ne11); + + const int64_t i01 = src1_dd[i10*s10 + i11*s11 + i12*s12]; + + float * y = dst_dd + i10*s1 + i11*s2 + i12*s3 + b*QK_K; + dequantize_block(y, i01*nblk + b, get_pointer(scales_local), item_ct1, nb); + }; + + if constexpr (dims == 1) { + cgh.parallel_for(sycl::nd_range<1>(sycl::range<1>(ngroups * nthreads), sycl::range<1>(nthreads)), kernel); + } else { + cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, ngroups * nthreads), sycl::range<3>(1, 1, nthreads)), kernel); + } + }); + + GGML_UNUSED(src0); +} + template static void get_rows_sycl_f32(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1, ggml_tensor *dst, const void *src0_dd, @@ -332,6 +460,9 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { GGML_ASSERT(dst->nb[0] == ggml_type_size(dst->type)); const int32_t * src1_i32 = (const int32_t *) dst->src[1]->data; + // a weight shared with MUL_MAT (e.g. tied embeddings) can be in reorder layout + const bool src0_reordered = dst->src[0]->extra && + ((const ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder; /* TODO: Refactor and remove duplicates */ switch (dst->src[0]->type) { case GGML_TYPE_F16: @@ -399,14 +530,37 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { src1_i32, (float *)dst->data, ctx.stream()); break; case GGML_TYPE_Q2_K: + if (src0_reordered) { + const void * vx = dst->src[0]->data; + get_rows_sycl_reorder_k<3, 64, 1>(dst->src[0], dst->src[1], dst, src1_i32, (float *)dst->data, ctx.stream(), + [=](float * y, int64_t ib, uint8_t * scales_local, const sycl::nd_item<3> & item_ct1, int64_t nb) { + dequantize_block_q2_K_reorder(vx, y, ib, item_ct1, nb); + GGML_UNUSED(scales_local); + }); + break; + } get_rows_sycl_f32(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, src1_i32, (float *)dst->data, ctx.stream()); break; case GGML_TYPE_Q3_K: + if (src0_reordered) { + const void * vx = dst->src[0]->data; + get_rows_sycl_reorder_k<3, 64, 1>(dst->src[0], dst->src[1], dst, src1_i32, (float *)dst->data, ctx.stream(), + [=](float * y, int64_t ib, uint8_t * scales_local, const sycl::nd_item<3> & item_ct1, int64_t nb) { + dequantize_block_q3_K_reorder(vx, y, ib, item_ct1, nb); + GGML_UNUSED(scales_local); + }); + break; + } get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, src1_i32, (float *)dst->data, ctx.stream()); break; case GGML_TYPE_Q4_0: + if (src0_reordered) { + get_rows_sycl_reorder(dst->src[0], dst->src[1], dst, dst->src[0]->data, + src1_i32, (float *)dst->data, ctx.stream()); + break; + } get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, src1_i32, (float *)dst->data, ctx.stream()); break; @@ -415,6 +569,14 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { src1_i32, (float *)dst->data, ctx.stream()); break; case GGML_TYPE_Q4_K: + if (src0_reordered) { + const void * vx = dst->src[0]->data; + get_rows_sycl_reorder_k<1, 32, K_SCALE_SIZE>(dst->src[0], dst->src[1], dst, src1_i32, (float *)dst->data, ctx.stream(), + [=](float * y, int64_t ib, uint8_t * scales_local, const sycl::nd_item<1> & item_ct1, int64_t nb) { + dequantize_block_q4_K_reorder(vx, y, ib, scales_local, item_ct1, nb); + }); + break; + } get_rows_sycl_f32(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, src1_i32, (float *)dst->data, ctx.stream()); break; @@ -427,14 +589,36 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { src1_i32, (float *)dst->data, ctx.stream()); break; case GGML_TYPE_Q5_K: + if (src0_reordered) { + const void * vx = dst->src[0]->data; + get_rows_sycl_reorder_k<3, 64, K_SCALE_SIZE>(dst->src[0], dst->src[1], dst, src1_i32, (float *)dst->data, ctx.stream(), + [=](float * y, int64_t ib, uint8_t * scales_local, const sycl::nd_item<3> & item_ct1, int64_t nb) { + dequantize_block_q5_K_reorder(vx, y, ib, scales_local, item_ct1, nb); + }); + break; + } get_rows_sycl_f32(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, src1_i32, (float *)dst->data, ctx.stream()); break; case GGML_TYPE_Q6_K: + if (src0_reordered) { + const void * vx = dst->src[0]->data; + get_rows_sycl_reorder_k<3, 64, 1>(dst->src[0], dst->src[1], dst, src1_i32, (float *)dst->data, ctx.stream(), + [=](float * y, int64_t ib, uint8_t * scales_local, const sycl::nd_item<3> & item_ct1, int64_t nb) { + dequantize_block_q6_K_reorder(vx, y, ib, item_ct1, nb); + GGML_UNUSED(scales_local); + }); + break; + } get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, src1_i32, (float *)dst->data, ctx.stream()); break; case GGML_TYPE_Q8_0: + if (src0_reordered) { + get_rows_sycl_reorder(dst->src[0], dst->src[1], dst, dst->src[0]->data, + src1_i32, (float *)dst->data, ctx.stream()); + break; + } get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, src1_i32, (float *)dst->data, ctx.stream()); break; diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 8fc90e037..1df1129cd 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -2453,6 +2453,63 @@ struct test_get_rows : public test_case { } }; +// GGML_OP_GET_ROWS on a weight also used by GGML_OP_MUL_MAT (tied embeddings) +// the mul_mat runs first, so a backend that changes the weight layout must still return the right rows +struct test_get_rows_after_mul_mat : public test_case { + const ggml_type type; + const int n; // cols + const int m; // rows + const int r; // rows to get + + std::string vars() override { + return VARS_TO_STR4(type, n, m, r); + } + + test_get_rows_after_mul_mat(ggml_type type = GGML_TYPE_Q8_0, int n = 256, int m = 16, int r = 4) + : type(type), n(n), m(m), r(r) {} + + double max_nmse_err() override { + return 5e-4; // same as test_mul_mat + } + + ggml_tensor * build_graph(ggml_context * ctx) override { + ggml_tensor * w = ggml_new_tensor_2d(ctx, type, n, m); + ggml_set_name(w, "w"); + + ggml_tensor * x = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, n, 1); + ggml_set_name(x, "x"); + + ggml_tensor * rows = ggml_new_tensor_1d(ctx, GGML_TYPE_I32, r); + ggml_set_name(rows, "rows"); + + ggml_tensor * mm = ggml_mul_mat(ctx, w, x); + ggml_set_name(mm, "mm"); + + ggml_tensor * gr = ggml_get_rows(ctx, w, rows); + ggml_set_name(gr, "gr"); + + // repeat(mm) is the first source, so mul_mat is computed before get_rows + ggml_tensor * out = ggml_add(ctx, ggml_repeat(ctx, mm, gr), gr); + ggml_set_name(out, "out"); + + return out; + } + + void initialize_tensors(ggml_context * ctx) override { + for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) { + if (t->type == GGML_TYPE_I32) { + std::vector data(ggml_nelements(t)); + for (size_t i = 0; i < data.size(); i++) { + data[i] = rand() % m; + } + ggml_backend_tensor_set(t, data.data(), 0, data.size() * sizeof(int)); + } else { + init_tensor_uniform(t); + } + } + } +}; + // GGML_OP_GET_ROWS_BACK struct test_get_rows_back : public test_case { const ggml_type type; @@ -9828,6 +9885,9 @@ static std::vector> make_test_cases_eval() { } test_cases.emplace_back(new test_get_rows(GGML_TYPE_F32, 1, 8, 2, 1, 1, false)); + for (ggml_type type : {GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, GGML_TYPE_Q4_K, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K}) { + test_cases.emplace_back(new test_get_rows_after_mul_mat(type)); + } for (ggml_type type : all_types) { for (int b : {1, 7}) { for (bool v : {false, true}) {