diff --git a/ggml/src/ggml-sycl/convert.cpp b/ggml/src/ggml-sycl/convert.cpp index 67dfeb9f9..8ba26d3a9 100644 --- a/ggml/src/ggml-sycl/convert.cpp +++ b/ggml/src/ggml-sycl/convert.cpp @@ -158,16 +158,16 @@ static void dequantize_row_q4_0_sycl_reorder(const void *vx, dst_t *y, const int dpct::has_capability_or_fail(stream->get_device(), {sycl::aspect::fp16}); - int constexpr WARP_K = WARP_SIZE * QK4_0; - const int n_warp = (k + WARP_K - 1) / WARP_K; - GGML_ASSERT(k % 2 == 0); - stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, n_warp) * - sycl::range<3>(1, 1, WARP_SIZE), - sycl::range<3>(1, 1, WARP_SIZE)), - [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{ + GGML_ASSERT(k % QK4_0 == 0); + const int64_t n_groups = ceil_div(k / 2, SYCL_DEQUANTIZE_BLOCK_SIZE); + // decrease global range when it exceeds the max int; the kernel loops to cover the rest + const int64_t wg_size = downsample_sycl_global_range(n_groups, SYCL_DEQUANTIZE_BLOCK_SIZE); + stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, n_groups) * + sycl::range<3>(1, 1, wg_size), + sycl::range<3>(1, 1, wg_size)), + [=](sycl::nd_item<3> item_ct1) { dequantize_block_q4_0_reorder(vx, y, k, item_ct1); }); - } template @@ -177,16 +177,16 @@ static void dequantize_row_q8_0_sycl_reorder(const void *vx, dst_t *y, const int dpct::has_capability_or_fail(stream->get_device(), {sycl::aspect::fp16}); - int constexpr WARP_K = WARP_SIZE * QK8_0; - const int n_warp = (k + WARP_K - 1) / WARP_K; GGML_ASSERT(k % QK8_0 == 0); - stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, n_warp) * - sycl::range<3>(1, 1, WARP_SIZE), - sycl::range<3>(1, 1, WARP_SIZE)), - [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{ + const int64_t n_groups = ceil_div(k / 2, SYCL_DEQUANTIZE_BLOCK_SIZE); + // decrease global range when it exceeds the max int; the kernel loops to cover the rest + const int64_t wg_size = downsample_sycl_global_range(n_groups, SYCL_DEQUANTIZE_BLOCK_SIZE); + stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, n_groups) * + sycl::range<3>(1, 1, wg_size), + sycl::range<3>(1, 1, wg_size)), + [=](sycl::nd_item<3> item_ct1) { dequantize_block_q8_0_reorder(vx, y, k, item_ct1); }); - } template diff --git a/ggml/src/ggml-sycl/dequantize.hpp b/ggml/src/ggml-sycl/dequantize.hpp index 14529a26b..48ee651ad 100644 --- a/ggml/src/ggml-sycl/dequantize.hpp +++ b/ggml/src/ggml-sycl/dequantize.hpp @@ -818,61 +818,43 @@ static void dequantize_block_q4_0(const void * __restrict__ vx, dst_t * __restri } template -static void dequantize_block_q4_0_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, int64_t nb32, +static void dequantize_block_q4_0_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, int64_t k, const sycl::nd_item<3> &item_ct1) { + // one thread per qs byte, like dequantize_block; the reorder layout makes the byte reads contiguous + // grid-stride loop because the launcher may shrink the work-group to keep the global range below INT_MAX + // float math, not dequantize_q4_0_reorder: that helper computes in half with GGML_SYCL_F16, which changes f32 output + const sycl::half * d_ptr = (const sycl::half *) ((const uint8_t *) vx + k/2); + const int64_t stride = item_ct1.get_local_range(2) * item_ct1.get_group_range(2); - const int64_t i = item_ct1.get_group(2); - auto k=nb32; - // assume 32 threads - const int64_t tid = item_ct1.get_local_id(2); - const int lane_ib = i * WARP_SIZE + tid; + for (int64_t iq = item_ct1.get_global_id(2); iq < k/2; iq += stride) { + const int64_t ib = iq / (QK4_0/2); + const int64_t iqs = iq % (QK4_0/2); + const int vq = ((const uint8_t *) vx)[iq]; + const float d = float(d_ptr[ib]); - if (lane_ib >= k / QK4_0) { - return; + dst_t * y = yy + ib*QK4_0 + iqs; + y[0] = d * ((vq & 0xF) - 8); + y[QK4_0/2] = d * ((vq >> 4) - 8); } - - dst_t * y_ptr = yy + lane_ib * QK4_0; - - auto qs = (const uint8_t*)vx + lane_ib * QK4_0 / 2; - auto s_ptr = (const sycl::half*)((const uint8_t*)vx + k / 2) + lane_ib; - - const float d = float(*s_ptr); - -#pragma unroll - for (int l = 0; l < QK4_0 / 2; ++l) { - int vq = qs[l]; - y_ptr[l + 0] = d * ((vq & 0xF) - 8); - y_ptr[l + 16] = d * ((vq >> 4) - 8); - } - } // Dequantize Q8_0 from reorder layout: [all qs (k bytes)][all d values] -// Each thread handles one block of QK8_0 elements. template static void dequantize_block_q8_0_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, int64_t k, const sycl::nd_item<3> &item_ct1) { + // two values per thread, like dequantize_block; the reorder layout keeps all qs contiguous + // grid-stride loop because the launcher may shrink the work-group to keep the global range below INT_MAX + // float math, not dequantize_q8_0_reorder: that helper computes in half with GGML_SYCL_F16, which changes f32 output + const int8_t * qs = (const int8_t *) vx; + const sycl::half * d_ptr = (const sycl::half *) ((const uint8_t *) vx + k); + const int64_t stride = 2 * item_ct1.get_local_range(2) * item_ct1.get_group_range(2); - const int64_t i = item_ct1.get_group(2); - const int64_t tid = item_ct1.get_local_id(2); - const int lane_ib = i * WARP_SIZE + tid; + for (int64_t i = 2 * item_ct1.get_global_id(2); i < k; i += stride) { + const float d = float(d_ptr[i / QK8_0]); - if (lane_ib >= k / QK8_0) { - return; + yy[i + 0] = d * qs[i + 0]; + yy[i + 1] = d * qs[i + 1]; } - - dst_t * y_ptr = yy + lane_ib * QK8_0; - - auto qs = (const int8_t*)vx + lane_ib * QK8_0; - auto s_ptr = (const sycl::half*)((const uint8_t*)vx + k) + lane_ib; - - const float d = float(*s_ptr); - -#pragma unroll - for (int l = 0; l < QK8_0; ++l) { - y_ptr[l] = d * qs[l]; - } - } template