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

fix buf

Browse source
This commit is contained in:
luoyu-intel 2024-07-16 06:37:10 +00:00
parent 216201230c
commit a8c75c041d
1 changed files with 6 additions and 3 deletions
Download patch file Download diff file Expand all files Collapse all files

9
ggml/src/ggml-sycl.cpp
View file

@ -4341,6 +4341,9 @@ static void ggml_backend_sycl_buffer_set_tensor(ggml_backend_buffer_t buffer,
if (tensor->type == GGML_TYPE_Q4_0) if (tensor->type == GGML_TYPE_Q4_0)
{ {
auto tmp_buf = sycl::malloc_shared<char>(size, *stream); auto tmp_buf = sycl::malloc_shared<char>(size, *stream);
SYCL_CHECK(
CHECK_TRY_ERROR((*stream).memcpy(tmp_buf, data, size)
.wait()));
GGML_ASSERT((size % sizeof(block_q4_0) == 0)); GGML_ASSERT((size % sizeof(block_q4_0) == 0));
GGML_ASSERT((offset % sizeof(block_q4_0) == 0)); GGML_ASSERT((offset % sizeof(block_q4_0) == 0));
int blk_offset = offset / sizeof(block_q4_0); int blk_offset = offset / sizeof(block_q4_0);
@ -4349,9 +4352,9 @@ static void ggml_backend_sycl_buffer_set_tensor(ggml_backend_buffer_t buffer,
size_t nrows = tensor->ne[1]; size_t nrows = tensor->ne[1];
auto d_ptr = (sycl::half*)(qs_ptr + ncols * nrows / 2) + blk_offset; auto d_ptr = (sycl::half*)(qs_ptr + ncols * nrows / 2) + blk_offset;
stream->parallel_for( stream->parallel_for(
size / QK4_0, size / sizeof(block_q4_0),
[=](auto i) [[intel::reqd_sub_group_size(WARP_SIZE)]] { [=](auto i) [[intel::reqd_sub_group_size(WARP_SIZE)]] {
const block_q4_0* x = (const block_q4_0*)data; const block_q4_0* x = (const block_q4_0*)tmp_buf;
int ib = i; int ib = i;
typedef sycl::vec<uint8_t, QK4_0 / 2> CT; typedef sycl::vec<uint8_t, QK4_0 / 2> CT;
CT tmp = *(CT*)x[ib].qs; CT tmp = *(CT*)x[ib].qs;
@ -4371,7 +4374,7 @@ static void ggml_backend_sycl_buffer_set_tensor(ggml_backend_buffer_t buffer,
} }
*(d_ptr + ib) = x[ib].d; *(d_ptr + ib) = x[ib].d;
}); }).wait();
sycl::free(tmp_buf, *stream); sycl::free(tmp_buf, *stream);
} }
else else