From f62dc45f318e48d375e7734b34cbddee81deed52 Mon Sep 17 00:00:00 2001
From: Akarshan Biswas <akarshan.biswas@gmail.com>
Date: Sun, 5 Jan 2025 18:11:38 +0530
Subject: [PATCH 1/3] SYCL: Use get_multi_ptr instead of deprecated get_pointer
 in wkv6

---
 ggml/src/ggml-sycl/wkv6.cpp | 12 ++++--------
 1 file changed, 4 insertions(+), 8 deletions(-)

diff --git a/ggml/src/ggml-sycl/wkv6.cpp b/ggml/src/ggml-sycl/wkv6.cpp
index 75ddfb86ac0f7..712090ddad25e 100644
--- a/ggml/src/ggml-sycl/wkv6.cpp
+++ b/ggml/src/ggml-sycl/wkv6.cpp
@@ -126,14 +126,10 @@ 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());
+        });
     });
 
     GGML_UNUSED(src0);

From 198fc8c9010de5c83ffc7af10169e06fb9f9c1c9 Mon Sep 17 00:00:00 2001
From: Akarshan Biswas <akarshan.biswas@gmail.com>
Date: Tue, 7 Jan 2025 11:15:51 +0530
Subject: [PATCH 2/3] 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 712090ddad25e..75ddfb86ac0f7 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<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,
-                                (float*)shared_mem_acc.get_multi_ptr<sycl::access::decorated::no>().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);

From 94081b6fd86adf9e71baf89863fdd8a0786cfa40 Mon Sep 17 00:00:00 2001
From: Akarshan Biswas <akarshan.biswas@gmail.com>
Date: Tue, 7 Jan 2025 11:17:37 +0530
Subject: [PATCH 3/3] Reland: Use get_multi_ptr instead of deprecated
 get_pointer in wkv6

---
 ggml/src/ggml-sycl/wkv6.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/ggml/src/ggml-sycl/wkv6.cpp b/ggml/src/ggml-sycl/wkv6.cpp
index 75ddfb86ac0f7..105db6f030c59 100644
--- a/ggml/src/ggml-sycl/wkv6.cpp
+++ b/ggml/src/ggml-sycl/wkv6.cpp
@@ -131,7 +131,7 @@ void ggml_sycl_op_rwkv_wkv6(ggml_backend_sycl_context& ctx, const ggml_tensor* s
             [=](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()
+                    item_ct1, (float*)shared_mem_acc.get_multi_ptr<sycl::access::decorated::no>().get()
                 );
             });
     });