vbatts/llama.cpp
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

fix deq kernel

Browse source
This commit is contained in:
luoyu-intel 2024-07-18 10:39:32 +08:00
parent e62fab8a49
commit 20d3b74020
1 changed files with 4 additions and 4 deletions
Download patch file Download diff file Expand all files Collapse all files

8
ggml/src/ggml-sycl/convert.cpp
View file

@ -125,7 +125,7 @@ static void dequantize_row_q4_0_sycl(const void *vx, dst_t *y, const int k,
const int warp_id = item_ct1.get_group(2); const int warp_id = item_ct1.get_group(2);
const int lane_id = item_ct1.get_local_id(2); const int lane_id = item_ct1.get_local_id(2);
const int lane_ib = warp_id * WARP_SIZE + lane_id; const int lane_ib = warp_id * WARP_SIZE + lane_id;
if (lane_ib >= k) { if (lane_ib >= k / Q4_0) {
return; return;
} }
@ -136,11 +136,11 @@ static void dequantize_row_q4_0_sycl(const void *vx, dst_t *y, const int k,
const float d = float(*s_ptr); const float d = float(*s_ptr);
auto qs = *(sycl::vec<uint8_t, QK4_0 / 2>*)q_ptr; auto qs = *(sycl::vec<uint8_t, QK4_0 / 2>*)q_ptr;
#pragma unroll
for (int l = 0; l < QK4_0 / 2; ++l) { for (int l = 0; l < QK4_0 / 2; ++l) {
int vq = qs[l]; int vq = qs[l];
y[l + 0] = d * ((vq & 0xF) - 8); y[l * 2 + 0] = d * ((vq & 0xF) - 8);
y[l + 16] = d * ((vq >> 4) - 8); y[l * 2 + 1] = d * ((vq >> 4) - 8);
} }
}); });
#else #else