Skip to content

Commit 198fc8c

Browse files
committed
Revert "SYCL: Use get_multi_ptr instead of deprecated get_pointer in wkv6"
This reverts commit f62dc45.
1 parent f62dc45 commit 198fc8c

File tree

1 file changed

+8
-4
lines changed

1 file changed

+8
-4
lines changed

ggml/src/ggml-sycl/wkv6.cpp

Lines changed: 8 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -126,10 +126,14 @@ void ggml_sycl_op_rwkv_wkv6(ggml_backend_sycl_context& ctx, const ggml_tensor* s
126126
stream->submit([&](sycl::handler& cgh) {
127127
sycl::local_accessor<float, 1> shared_mem_acc(shared_mem_size, cgh);
128128

129-
cgh.parallel_for(sycl::nd_range<3>(grid_dims * block_dims, block_dims), [=](sycl::nd_item<3> item_ct1) {
130-
rwkv_wkv_f32_kernel(B, T, C, H, k_d, v_d, r_d, tf_d, td_d, s_d, dst_d, item_ct1,
131-
(float*)shared_mem_acc.get_multi_ptr<sycl::access::decorated::no>().get());
132-
});
129+
cgh.parallel_for(
130+
sycl::nd_range<3>(grid_dims * block_dims, block_dims),
131+
[=](sycl::nd_item<3> item_ct1) {
132+
rwkv_wkv_f32_kernel(
133+
B, T, C, H, k_d, v_d, r_d, tf_d, td_d, s_d, dst_d,
134+
item_ct1, shared_mem_acc.get_pointer()
135+
);
136+
});
133137
});
134138

135139
GGML_UNUSED(src0);

0 commit comments

Comments
 (0)