diff --git a/app_neso_advection_mdir_cg/hip/move_kernel_loop.hpp b/app_neso_advection_mdir_cg/hip/move_kernel_loop.hpp
index d7df231..8c4130b 100644
--- a/app_neso_advection_mdir_cg/hip/move_kernel_loop.hpp
+++ b/app_neso_advection_mdir_cg/hip/move_kernel_loop.hpp
@@ -237,7 +237,7 @@ void opp_particle_move__move_kernel(opp_set set, opp_map c2c_map, opp_map p2c_ma
  
     opp_profiler->end("move_kernel");
 }
-void opp_init_direct_hop_cg(double grid_spacing, int dim, const opp_dat c_gbl_id, const opp::BoundingBox& b_box, 
+void opp_init_direct_hop_cg(double grid_spacing, const opp_dat c_gbl_id, const opp::BoundingBox& b_box, 
     opp_map c2c_map, opp_map p2c_map,
     opp_arg arg0, // p_pos | OPP_READ
     opp_arg arg1, // p_mdir | OPP_RW
diff --git a/app_neso_advection_mdir_cg/sycl/move_kernel_loop.hpp b/app_neso_advection_mdir_cg/sycl/move_kernel_loop.hpp
index 41b246e..f9c42d6 100644
--- a/app_neso_advection_mdir_cg/sycl/move_kernel_loop.hpp
+++ b/app_neso_advection_mdir_cg/sycl/move_kernel_loop.hpp
@@ -207,7 +207,7 @@ void opp_particle_move__move_kernel(opp_set set, opp_map c2c_map, opp_map p2c_ma
  
     opp_profiler->end("move_kernel");
 }
-void opp_init_direct_hop_cg(double grid_spacing, int dim, const opp_dat c_gbl_id, const opp::BoundingBox& b_box, 
+void opp_init_direct_hop_cg(double grid_spacing, const opp_dat c_gbl_id, const opp::BoundingBox& b_box, 
     opp_map c2c_map, opp_map p2c_map,
     opp_arg arg0, // p_pos | OPP_READ
     opp_arg arg1, // p_mdir | OPP_RW
diff --git a/opp_lib/src/sycl/opp_direct_hop_sycl.cpp b/opp_lib/src/sycl/opp_direct_hop_sycl.cpp
index e5cd150..2455de0 100644
--- a/opp_lib/src/sycl/opp_direct_hop_sycl.cpp
+++ b/opp_lib/src/sycl/opp_direct_hop_sycl.cpp
@@ -870,7 +870,7 @@ void dh_particle_packer_gpu::pack(opp_set set)
     opp_profiler->start("MvDH_Pack");
 
     std::map<int, std::vector<char>>& buffers_of_set = this->buffers[set->index];
-    thrust::device_vector<OPP_INT>& temp_dv = *(set->mesh_relation_dat->thrust_int_sort);
+    OPP_INT* temp_dp = (OPP_INT*)set->mesh_relation_dat->data_swap_d;
 
     for (const auto& per_rank_parts : local_part_ids) {
 
@@ -887,7 +887,6 @@ void dh_particle_packer_gpu::pack(opp_set set)
 
         const int copy_count = (int)part_ids_vec.size();
         
-        OPP_INT* temp_dp = opp_get_dev_raw_ptr<OPP_INT>(temp_dv);
         opp_mem::copy_host_to_dev<OPP_INT>(temp_dp, part_ids_vec.data(), copy_count);
 
         size_t disp = 0;
@@ -900,19 +899,19 @@ void dh_particle_packer_gpu::pack(opp_set set)
             }
             else if (strcmp(dat->type, "double") == 0) {
                     
-                copy_according_to_index<OPP_REAL>(dat->thrust_real, dat->thrust_real_sort, 
-                    temp_dv, dat->set->set_capacity, copy_count, copy_count, dat->dim);
+                copy_according_to_index<OPP_REAL>((OPP_REAL*)dat->data_d, (OPP_REAL*)dat->data_swap_d,
+                    temp_dp, dat->set->set_capacity, copy_count, copy_count, dat->dim);
 
                 opp_mem::copy_dev_to_host<char>(&(send_rank_buffer[disp]), 
-                    (char*)opp_get_dev_raw_ptr<OPP_REAL>(*(dat->thrust_real_sort)), bytes_to_copy);
+                    dat->data_swap_d, bytes_to_copy);
             }
             else if (strcmp(dat->type, "int") == 0) {
                 
-                copy_according_to_index<OPP_INT>(dat->thrust_int, dat->thrust_int_sort, 
-                    temp_dv, dat->set->set_capacity, copy_count, copy_count, dat->dim);
+                copy_according_to_index<OPP_INT>((OPP_INT*)dat->data_d, (OPP_INT*)dat->data_swap_d, 
+                    temp_dp, dat->set->set_capacity, copy_count, copy_count, dat->dim);
                 
                 opp_mem::copy_dev_to_host<char>(&(send_rank_buffer[disp]), 
-                    (char*)opp_get_dev_raw_ptr<OPP_INT>(*(dat->thrust_int_sort)), bytes_to_copy);
+                    dat->data_swap_d, bytes_to_copy);
             }
 
             disp += bytes_to_copy; 
diff --git a/opp_translator/resources/templates/cpp/hip/move_loop_host.hpp.jinja b/opp_translator/resources/templates/cpp/hip/move_loop_host.hpp.jinja
index 14ab332..acece1d 100644
--- a/opp_translator/resources/templates/cpp/hip/move_loop_host.hpp.jinja
+++ b/opp_translator/resources/templates/cpp/hip/move_loop_host.hpp.jinja
@@ -873,7 +873,7 @@ opp_dev_sr_{{lh.kernel}}<<<num_blocks, block_size{{-(", %s" % shared_size) if lh
 
 {% block dh_init_wrapper %}
 {% if lh.dh_loop_required %}
-void opp_init_direct_hop_cg(double grid_spacing, int dim, const opp_dat c_gbl_id, const opp::BoundingBox& b_box, 
+void opp_init_direct_hop_cg(double grid_spacing, const opp_dat c_gbl_id, const opp::BoundingBox& b_box, 
     opp_map c2c_map, opp_map p2c_map,
     {% for arg in lh.args %}
     opp_arg arg{{arg.id}}{{"," if not loop.last}} // {% if arg is dat %}{{lh.dat(arg).ptr}} {% endif -%} | OPP_{{arg.access_type.name}}
diff --git a/opp_translator/resources/templates/cpp/sycl/move_loop_host.hpp.jinja b/opp_translator/resources/templates/cpp/sycl/move_loop_host.hpp.jinja
index 12f8a37..2f06d74 100644
--- a/opp_translator/resources/templates/cpp/sycl/move_loop_host.hpp.jinja
+++ b/opp_translator/resources/templates/cpp/sycl/move_loop_host.hpp.jinja
@@ -474,7 +474,7 @@ auto opp_part_check_status =
 
 {% block dh_init_wrapper %}
 {% if lh.dh_loop_required %}
-void opp_init_direct_hop_cg(double grid_spacing, int dim, const opp_dat c_gbl_id, const opp::BoundingBox& b_box, 
+void opp_init_direct_hop_cg(double grid_spacing, const opp_dat c_gbl_id, const opp::BoundingBox& b_box, 
     opp_map c2c_map, opp_map p2c_map,
     {% for arg in lh.args %}
     opp_arg arg{{arg.id}}{{"," if not loop.last}} // {% if arg is dat %}{{lh.dat(arg).ptr}} {% endif -%} | OPP_{{arg.access_type.name}}