passing the accessors with their ptrs

ginkgo-project · Mar 9, 2023 · d24a473 · d24a473
1 parent 997ef8c
commit d24a473
Show file tree

Hide file tree

Showing 3 changed files with 11 additions and 15 deletions.
diff --git 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
@@ -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
@@ -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);