darthrubik/llama-cpp
1
0
Fork 0
Code Issues Pull Requests 1 Packages Projects Releases Wiki Activity

sycl: q8 fast dequant (#30775)
Copilot Setup Steps / copilot-setup-steps (push) Has been cancelled
Check Pre-Tokenizer Hashes / pre-tokenizer-hashes (push) Has been cancelled
Python check requirements.txt / check-requirements (push) Has been cancelled
Python Type-Check / python type-check (push) Has been cancelled
Update Operations Documentation / update-ops-docs (push) Has been cancelled

Browse Source
This commit is contained in:
92es
2026-10-21 10:01:03 -07:00
committed by GitHub
parent 9febe63d5a
commit 0756e83fc1
2 changed files with 25 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-sycl/convert.cpp
+5 -8
View File
@@ -177,15 +177,12 @@ 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});
// flat range<1>, two pairs (4 elements) per item: no nd_range/group math, one vectorized store.
// k % QK8_0 == 0 implies k % 4 == 0, so the wide launch covers the row exactly (no tail).
GGML_ASSERT(k % QK8_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_q8_0_reorder(vx, y, k, item_ct1);
stream->parallel_for(sycl::range<1>((size_t) (k / 4)),
[=](sycl::id<1> idx) {
dequantize_block_q8_0_reorder_wide(vx, y, k, idx);
});
}
ggml/src/ggml-sycl/dequantize.hpp
+20
View File
@@ -857,6 +857,26 @@ static void dequantize_block_q8_0_reorder(const void * __restrict__ vx, dst_t *
}
}
template<typename dst_t>
static void dequantize_block_q8_0_reorder_wide(const void * __restrict__ vx, dst_t * __restrict__ yy, int64_t k,
const sycl::id<1> idx) {
const int64_t i = 4 * (int64_t) idx[0];
const int8_t * qs = (const int8_t *) vx;
const sycl::half * d_ptr = (const sycl::half *) ((const uint8_t *) vx + k);
const float d0 = float(d_ptr[i / QK8_0]);
const float d1 = float(d_ptr[(i + 2) / QK8_0]);
sycl::vec<float, 4> p;
p[0] = d0 * (float) qs[i + 0];
p[1] = d0 * (float) qs[i + 1];
p[2] = d1 * (float) qs[i + 2];
p[3] = d1 * (float) qs[i + 3];
*reinterpret_cast<sycl::vec<dst_t, 4> *>(yy + i) = p.convert<dst_t, sycl::rounding_mode::automatic>();
}
template<typename dst_t>
static void dequantize_block_q4_1(const void * __restrict__ vx, dst_t * __restrict__ yy, int64_t nb32,
const sycl::nd_item<3> &item_ct1) {