From 4a481556e606ecd06eaf7da628aa830874130d68 Mon Sep 17 00:00:00 2001 From: Aidan Date: Mon, 17 Jun 2024 10:16:10 +0100 Subject: [PATCH 1/4] Remove double lines --- ggml-sycl.cpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/ggml-sycl.cpp b/ggml-sycl.cpp index 4702c788b9e8d..118baa268d249 100644 --- a/ggml-sycl.cpp +++ b/ggml-sycl.cpp @@ -4297,7 +4297,8 @@ static void dequantize_block_q3_K(const void * __restrict__ vx, dst_t * __restri static inline void get_scale_min_k4(int j, const uint8_t * q, uint8_t & d, uint8_t & m) { if (j < 4) { - d = q[j] & 63; m = q[j + 4] & 63; + d = q[j] & 63; + m = q[j + 4] & 63; } else { d = (q[j+4] & 0xF) | ((q[j-4] >> 6) << 4); m = (q[j+4] >> 4) | ((q[j-0] >> 6) << 4); @@ -4327,9 +4328,11 @@ static void dequantize_block_q4_K(const void * __restrict__ vx, dst_t * __restri uint8_t sc, m; get_scale_min_k4(is + 0, x[i].scales, sc, m); - const float d1 = dall * sc; const float m1 = dmin * m; + const float d1 = dall * sc; + const float m1 = dmin * m; get_scale_min_k4(is + 1, x[i].scales, sc, m); - const float d2 = dall * sc; const float m2 = dmin * m; + const float d2 = dall * sc; + const float m2 = dmin * m; for (int l = 0; l < n; ++l) { y[l + 0] = d1 * (q[l] & 0xF) - m1; y[l +32] = d2 * (q[l] >> 4) - m2; From cb3fb42046916cabe2e10482c1041251a79b7afb Mon Sep 17 00:00:00 2001 From: Aidan Date: Mon, 17 Jun 2024 10:21:16 +0100 Subject: [PATCH 2/4] Single load for half2 --- ggml-sycl.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/ggml-sycl.cpp b/ggml-sycl.cpp index 118baa268d249..1dc117af2d707 100644 --- a/ggml-sycl.cpp +++ b/ggml-sycl.cpp @@ -4321,8 +4321,9 @@ static void dequantize_block_q4_K(const void * __restrict__ vx, dst_t * __restri dst_t * y = yy + i*QK_K + 64*il + n*ir; - const float dall = x[i].dm[0]; - const float dmin = x[i].dm[1]; + const sycl::half2 dm = x[i].dm; + const float dall = dm[0]; + const float dmin = dm[1]; const uint8_t * q = x[i].qs + 32*il + n*ir; From 604ef6bf15abe28b24b4bccd59c3cb02739bdce4 Mon Sep 17 00:00:00 2001 From: Aidan Date: Mon, 17 Jun 2024 10:26:18 +0100 Subject: [PATCH 3/4] Store scales in local mem --- ggml-sycl.cpp | 17 ++++++++++++----- 1 file changed, 12 insertions(+), 5 deletions(-) diff --git a/ggml-sycl.cpp b/ggml-sycl.cpp index 1dc117af2d707..7b48c41bd5616 100644 --- a/ggml-sycl.cpp +++ b/ggml-sycl.cpp @@ -4307,7 +4307,7 @@ static inline void get_scale_min_k4(int j, const uint8_t * q, uint8_t & d, uint8 template static void dequantize_block_q4_K(const void * __restrict__ vx, dst_t * __restrict__ yy, - const sycl::nd_item<3> &item_ct1) { + uint8_t* scales_local, const sycl::nd_item<3> &item_ct1) { const block_q4_K * x = (const block_q4_K *) vx; const int i = item_ct1.get_group(2); @@ -4325,13 +4325,17 @@ static void dequantize_block_q4_K(const void * __restrict__ vx, dst_t * __restri const float dall = dm[0]; const float dmin = dm[1]; + if (tid < 12) + scales_local[tid] = x[i].scales[tid]; + item_ct1.barrier(sycl::access::fence_space::local_space); + const uint8_t * q = x[i].qs + 32*il + n*ir; uint8_t sc, m; - get_scale_min_k4(is + 0, x[i].scales, sc, m); + get_scale_min_k4(is + 0, scales_local, sc, m); const float d1 = dall * sc; const float m1 = dmin * m; - get_scale_min_k4(is + 1, x[i].scales, sc, m); + get_scale_min_k4(is + 1, scales_local, sc, m); const float d2 = dall * sc; const float m2 = dmin * m; for (int l = 0; l < n; ++l) { @@ -9894,12 +9898,15 @@ static void dequantize_row_q4_K_sycl(const void *vx, dst_t *y, const int k, dpct::has_capability_or_fail(stream->get_device(), {sycl::aspect::fp16}); - stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * + stream->submit([&](sycl::handler &cgh) { + sycl::local_accessor scale_local_acc(sycl::range<1>(12), cgh); + cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), [=](sycl::nd_item<3> item_ct1) { - dequantize_block_q4_K(vx, y, item_ct1); + dequantize_block_q4_K(vx, y, scale_local_acc.get_pointer(), item_ct1); }); + }); } } From a235b7c532bba6e5fd2f6d0c7215145800ed7c01 Mon Sep 17 00:00:00 2001 From: Aidan Date: Mon, 17 Jun 2024 10:30:40 +0100 Subject: [PATCH 4/4] Vectorize q load --- ggml-sycl.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/ggml-sycl.cpp b/ggml-sycl.cpp index 7b48c41bd5616..9d1956c180e69 100644 --- a/ggml-sycl.cpp +++ b/ggml-sycl.cpp @@ -4329,8 +4329,6 @@ static void dequantize_block_q4_K(const void * __restrict__ vx, dst_t * __restri scales_local[tid] = x[i].scales[tid]; item_ct1.barrier(sycl::access::fence_space::local_space); - const uint8_t * q = x[i].qs + 32*il + n*ir; - uint8_t sc, m; get_scale_min_k4(is + 0, scales_local, sc, m); const float d1 = dall * sc; @@ -4338,9 +4336,11 @@ static void dequantize_block_q4_K(const void * __restrict__ vx, dst_t * __restri get_scale_min_k4(is + 1, scales_local, sc, m); const float d2 = dall * sc; const float m2 = dmin * m; + + sycl::vec q_vec = reinterpret_cast*>(x[i].qs + 32*il + n*ir)[0]; for (int l = 0; l < n; ++l) { - y[l + 0] = d1 * (q[l] & 0xF) - m1; - y[l +32] = d2 * (q[l] >> 4) - m2; + y[l + 0] = d1 * (q_vec[l] & 0xF) - m1; + y[l +32] = d2 * (q_vec[l] >> 4) - m2; } }