vrr/koboldcpp
2
0
Fork 0
mirror of https://github.com/LostRuins/koboldcpp.git synced 2026-05-22 19:47:49 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

sycl: scalar SWAR byte-subtract in Q6_K MMVQ dot product (#22156)

Browse source 
Signed-off-by: Chun Tao <chun.tao@intel.com>
Co-authored-by: Chun Tao <chun.tao@intel.com>
This commit is contained in:
Intel AI Get-to Market Customer Success and Solutions 2026-05-17 22:12:21 -07:00 • committed by GitHub
parent 5511965b19
commit 0caf2a1d48
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
1 changed files with 46 additions and 53 deletions
Download patch file Download diff file Expand all files Collapse all files

99
ggml/src/ggml-sycl/vecdotq.hpp
View file

@ -85,6 +85,32 @@ static __dpct_inline__ int get_int_from_uint8_aligned(
(const int*)(x8 + sizeof(int) * i32)); // assume at least 4 byte alignment
}
static __dpct_inline__ int byte_sub_4(const int a, const int b) {
const uint32_t ua = static_cast<uint32_t>(a);
const uint32_t ub = static_cast<uint32_t>(b);
return static_cast<int>(((ua | 0x80808080u) - ub) ^ 0x80808080u);
}
static __dpct_inline__ float vec_dot_q6_K_q8_1_impl_mmvq_scalar(
const int vl, const int vh, const int u0, const int u1, const int8_t sc0,
const int8_t sc1, const float d, const float d80, const float d81) {
static_assert(QR6_K == 2, "q6_K MMVQ scalar fast path assumes QR6_K == 2");
const int vil0 = (vl >> 0) & 0x0F0F0F0F;
const int vih0 = ((vh >> 0) << 4) & 0x30303030;
const int vi0 = byte_sub_4(vil0 | vih0, 0x20202020);
const int vil1 = (vl >> 4) & 0x0F0F0F0F;
const int vih1 = ((vh >> 4) << 4) & 0x30303030;
const int vi1 = byte_sub_4(vil1 | vih1, 0x20202020);
const float sumf =
d80 * (dpct::dp4a(vi0, u0, 0) * sc0) +
d81 * (dpct::dp4a(vi1, u1, 0) * sc1);
return d * sumf;
}
static __dpct_inline__ void get_int_from_table_16(const uint32_t &q4,
const uint8_t *values,
int &val1, int &val2) {
@ -279,24 +305,8 @@ vec_dot_q6_K_q8_1_impl_mmvq(const int &vl, const int &vh,
const int *__restrict__ u,
const int8_t *__restrict__ scales, const float &d,
const float *__restrict__ d8) {
float sumf = 0.0f;
#pragma unroll
for (int i = 0; i < QR6_K; ++i) {
const int sc = scales[4*i];
const int vil = (vl >> (4*i)) & 0x0F0F0F0F;
const int vih = ((vh >> (4*i)) << 4) & 0x30303030;
const int vi = dpct::vectorized_binary<sycl::char4>(
(vil | vih), 0x20202020, dpct::sub_sat()); // vi = (vil | vih) - 32
sumf += d8[i] * (dpct::dp4a(vi, u[i], 0) * sc); // SIMD dot product
}
return d*sumf;
return vec_dot_q6_K_q8_1_impl_mmvq_scalar(
vl, vh, u[0], u[1], scales[0], scales[4], d, d8[0], d8[1]);
}
// VDR = vec dot ratio, how many contiguous integers each thread processes when the vec dot kernel is called
@ -542,23 +552,8 @@ template <> struct reorder_vec_dot_q_sycl<GGML_TYPE_Q6_K> {
__dpct_inline__ float vec_dot_q6_K_q8_1_impl_mmvq(const int vl, const int vh, const int * __restrict__ u,
const int8_t * __restrict__ scales, const float d,
const float * __restrict__ d8) {
float sumf = 0.0f;
#pragma unroll
for (int i = 0; i < QR6_K; ++i) {
const int sc = scales[4 * i];
const int vil = (vl >> (4 * i)) & 0x0F0F0F0F;
const int vih = ((vh >> (4 * i)) << 4) & 0x30303030;
const int vi = dpct::vectorized_binary<sycl::char4>((vil | vih), 0x20202020,
dpct::sub_sat()); // vi = (vil | vih) - 32
sumf += d8[i] * (dpct::dp4a(vi, u[i], 0) * sc); // SIMD dot product
}
return d * sumf;
return vec_dot_q6_K_q8_1_impl_mmvq_scalar(
vl, vh, u[0], u[1], scales[0], scales[4], d, d8[0], d8[1]);
}
__dpct_inline__ float operator()(const void * __restrict__ vbq, const std::pair<int, int> ibx_offset,
@ -579,16 +574,15 @@ template <> struct reorder_vec_dot_q_sycl<GGML_TYPE_Q6_K> {
const int8_t * scs = scales + scale_offset;
int u[QR6_K];
float d8[QR6_K];
const int u0 = get_int_from_int8_aligned(
q8_1_quant_ptr + bq8_offset * QK8_1, iqs % QI8_1);
const int u1 = get_int_from_int8_aligned(
q8_1_quant_ptr + (bq8_offset + 2) * QK8_1, iqs % QI8_1);
const float d80 = (*(q8_1_ds + bq8_offset + 0))[0];
const float d81 = (*(q8_1_ds + bq8_offset + 2))[0];
#pragma unroll
for (int i = 0; i < QR6_K; ++i) {
u[i] = get_int_from_int8_aligned(q8_1_quant_ptr + (bq8_offset + 2 * i) * QK8_1, iqs % QI8_1);
const sycl::half2 ds_values = *(q8_1_ds + bq8_offset + 2 * i);
d8[i] = ds_values[0];
}
return vec_dot_q6_K_q8_1_impl_mmvq(vl, vh, u, scs, *d, d8);
return vec_dot_q6_K_q8_1_impl_mmvq_scalar(
vl, vh, u0, u1, scs[0], scs[4], *d, d80, d81);
}
};
#define VDR_Q4_0_Q8_1_MMVQ 2
@ -1167,16 +1161,15 @@ vec_dot_q6_K_q8_1(const void *__restrict__ vbq,
const int8_t * scales = bq6_K->scales + scale_offset;
int u[QR6_K];
float d8[QR6_K];
const int u0 = get_int_from_int8_aligned(
bq8_1[bq8_offset + 0].qs, iqs % QI8_1);
const int u1 = get_int_from_int8_aligned(
bq8_1[bq8_offset + 2].qs, iqs % QI8_1);
const float d80 = bq8_1[bq8_offset + 0].ds[0];
const float d81 = bq8_1[bq8_offset + 2].ds[0];
#pragma unroll
for (int i = 0; i < QR6_K; ++i) {
u[i] = get_int_from_int8_aligned(bq8_1[bq8_offset + 2*i].qs, iqs % QI8_1);
d8[i] = bq8_1[bq8_offset + 2 * i].ds[0];
}
return vec_dot_q6_K_q8_1_impl_mmvq(vl, vh, u, scales, bq6_K->d, d8);
return vec_dot_q6_K_q8_1_impl_mmvq_scalar(
vl, vh, u0, u1, scales[0], scales[4], bq6_K->d, d80, d81);
}