root/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggerganov/llama.cpp.git synced 2025-01-12 03:31:46 +00:00
Code Issues Actions 5 Packages Projects Releases Wiki Activity

SYCL: Use get_multi_ptr instead of deprecated get_pointer in wkv6

Browse Source
This commit is contained in:
Akarshan Biswas 2025-01-05 18:11:38 +05:30
parent b56f079e28
commit f62dc45f31
No known key found for this signature in database
GPG Key ID: 52A578A14B32134D
1 changed files with 4 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
ggml/src/ggml-sycl/wkv6.cpp
View File

@ -126,13 +126,9 @@ void ggml_sycl_op_rwkv_wkv6(ggml_backend_sycl_context& ctx, const ggml_tensor* s
stream->submit([&](sycl::handler& cgh) {
sycl::local_accessor<float, 1> shared_mem_acc(shared_mem_size, cgh);
cgh.parallel_for(
sycl::nd_range<3>(grid_dims * block_dims, block_dims),
[=](sycl::nd_item<3> item_ct1) {
rwkv_wkv_f32_kernel(
B, T, C, H, k_d, v_d, r_d, tf_d, td_d, s_d, dst_d,
item_ct1, shared_mem_acc.get_pointer()
);
cgh.parallel_for(sycl::nd_range<3>(grid_dims * block_dims, block_dims), [=](sycl::nd_item<3> item_ct1) {
rwkv_wkv_f32_kernel(B, T, C, H, k_d, v_d, r_d, tf_d, td_d, s_d, dst_d, item_ct1,
(float*)shared_mem_acc.get_multi_ptr<sycl::access::decorated::no>().get());
});
});