working mdir advection sycl in feature/gpu_dh

OP-DSL · Dec 14, 2024 · 8bc450d · 8bc450d
1 parent 2ce4584
commit 8bc450d
Show file tree

Hide file tree

Showing 5 changed files with 11 additions and 12 deletions.
diff --git 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
@@ -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
@@ -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
@@ -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
@@ -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}}