From 198fc8c9010de5c83ffc7af10169e06fb9f9c1c9 Mon Sep 17 00:00:00 2001 From: Akarshan Biswas Date: Tue, 7 Jan 2025 11:15:51 +0530 Subject: [PATCH] Revert "SYCL: Use get_multi_ptr instead of deprecated get_pointer in wkv6" This reverts commit f62dc45f318e48d375e7734b34cbddee81deed52. --- ggml/src/ggml-sycl/wkv6.cpp | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-sycl/wkv6.cpp b/ggml/src/ggml-sycl/wkv6.cpp index 712090dda..75ddfb86a 100644 --- a/ggml/src/ggml-sycl/wkv6.cpp +++ b/ggml/src/ggml-sycl/wkv6.cpp @@ -126,10 +126,14 @@ void ggml_sycl_op_rwkv_wkv6(ggml_backend_sycl_context& ctx, const ggml_tensor* s stream->submit([&](sycl::handler& cgh) { sycl::local_accessor 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, - (float*)shared_mem_acc.get_multi_ptr().get()); - }); + 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() + ); + }); }); GGML_UNUSED(src0);