From d24a4738ba7854b6c006a15ded462572356efe32 Mon Sep 17 00:00:00 2001
From: Phuong Nguyen <bichphuong1204@gmail.com>
Date: Thu, 9 Mar 2023 17:07:24 +0000
Subject: [PATCH] passing the accessors with their ptrs

---
 dpcpp/preconditioner/batch_block_jacobi.hpp.inc  |  5 +----
 dpcpp/preconditioner/batch_jacobi_kernels.dp.cpp | 15 +++++++++------
 dpcpp/preconditioner/batch_scalar_jacobi.hpp.inc |  6 +-----
 3 files changed, 11 insertions(+), 15 deletions(-)

diff --git a/dpcpp/preconditioner/batch_block_jacobi.hpp.inc b/dpcpp/preconditioner/batch_block_jacobi.hpp.inc
index 9487df86c71..a18dcf2fa2f 100644
--- a/dpcpp/preconditioner/batch_block_jacobi.hpp.inc
+++ b/dpcpp/preconditioner/batch_block_jacobi.hpp.inc
@@ -164,10 +164,7 @@ template <typename ValueType>
 __dpct_inline__ void batch_block_jacobi_apply(
     BatchBlockJacobi<ValueType> prec, const size_type batch_id, const int nrows,
     const ValueType* const b_values, ValueType* const x_values,
-    sycl::accessor<ValueType, 1, sycl::access_mode::read_write,
-                   sycl::access::target::local>
-        sh_mem,
-    sycl::nd_item<3> item_ct1)
+    ValueType* sh_mem, sycl::nd_item<3> item_ct1)
 {
     ValueType* work = &sh_mem[0];
 
diff --git a/dpcpp/preconditioner/batch_jacobi_kernels.dp.cpp b/dpcpp/preconditioner/batch_jacobi_kernels.dp.cpp
index 4e403813ae4..06602f7f1e2 100644
--- a/dpcpp/preconditioner/batch_jacobi_kernels.dp.cpp
+++ b/dpcpp/preconditioner/batch_jacobi_kernels.dp.cpp
@@ -110,9 +110,11 @@ void batch_jacobi_apply_helper(
             cgh.parallel_for(
                 sycl_nd_range(grid, block), [=](sycl::nd_item<3> item_ct1) {
                     auto batch_id = item_ct1.get_group_linear_id();
-                    batch_scalar_jacobi_apply(prec_scalar_jacobi, sys_mat_batch,
-                                              batch_id, nrows, r_values,
-                                              z_values, slm_storage, item_ct1);
+                    batch_scalar_jacobi_apply(
+                        prec_scalar_jacobi, sys_mat_batch, batch_id, nrows,
+                        r_values, z_values,
+                        static_cast<ValueType*>(slm_storage.get_pointer()),
+                        item_ct1);
                 });
         });
 
@@ -133,9 +135,10 @@ void batch_jacobi_apply_helper(
             cgh.parallel_for(
                 sycl_nd_range(grid, block), [=](sycl::nd_item<3> item_ct1) {
                     auto batch_id = item_ct1.get_group_linear_id();
-                    batch_block_jacobi_apply(prec_block_jacobi, batch_id, nrows,
-                                             r_values, z_values, slm_storage,
-                                             item_ct1);
+                    batch_block_jacobi_apply(
+                        prec_block_jacobi, batch_id, nrows, r_values, z_values,
+                        static_cast<ValueType*>(slm_storage.get_pointer()),
+                        item_ct1);
                 });
         });
     }
diff --git a/dpcpp/preconditioner/batch_scalar_jacobi.hpp.inc b/dpcpp/preconditioner/batch_scalar_jacobi.hpp.inc
index d2f5207d6e4..311a54bb831 100644
--- a/dpcpp/preconditioner/batch_scalar_jacobi.hpp.inc
+++ b/dpcpp/preconditioner/batch_scalar_jacobi.hpp.inc
@@ -165,11 +165,7 @@ template <typename BatchMatrixType, typename ValueType>
 __dpct_inline__ void batch_scalar_jacobi_apply(
     BatchScalarJacobi<ValueType> prec, const BatchMatrixType sys_mat_batch,
     const size_type batch_id, const int nrows, const ValueType* const b_values,
-    ValueType* const x_values,
-    sycl::accessor<ValueType, 1, sycl::access_mode::read_write,
-                   sycl::access::target::local>
-        sh_mem,
-    sycl::nd_item<3> item_ct1)
+    ValueType* const x_values, ValueType* sh_mem, sycl::nd_item<3> item_ct1)
 {
     const auto sys_mat_batch_entry =
         batch::batch_entry(sys_mat_batch, batch_id);