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 __dpct_inline__ void batch_block_jacobi_apply( BatchBlockJacobi prec, const size_type batch_id, const int nrows, const ValueType* const b_values, ValueType* const x_values, - sycl::accessor - 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(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(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 __dpct_inline__ void batch_scalar_jacobi_apply( BatchScalarJacobi 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 - 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);