Compile hip and sycl backends on gpu_dh branch

app_fempic_cg/hip/move_kernel_loop.hpp

@@ -254,7 +254,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_lc | OPP_WRITE

app_neso_advection_cg/Makefile

@@ -149,25 +149,25 @@ sycl_mpi: mklib advec_opp+sycl_mpi opp_kernels+sycl_mpi
 
 # ------------------------------------------------------------------------------------------
 advec_opp+hip:
-	 $(HIPCC) $(HIPCCFLAGS) -DUSE_HIP -c advec_opp.cpp -o $(OBJ)/advec_opp_hip.o $(ALL_INC) $(HIP_INC)
+	 $(MPICPP) $(CPPFLAGS) -DUSE_HIP -c advec_opp.cpp -o $(OBJ)/advec_opp_hip.o $(ALL_INC) $(HIP_INC)
 opp_kernels+hip:
 	 $(HIPCC) $(HIPCCFLAGS) -DUSE_HIP -c hip/opp_kernels.cpp -o $(OBJ)/opp_kernels_hip.o $(ALL_INC) $(HIP_INC)
 hip: mklib advec_opp+hip opp_kernels+hip
-	 $(HIPCC) $(HIPCCFLAGS) -o $(BIN)/hip \
+	 $(MPICPP) $(CPPFLAGS) -o $(BIN)/hip \
 		$(OBJ)/advec_opp_hip.o \
 		$(OBJ)/opp_kernels_hip.o \
-		$(ALL_INC) $(HIP_INC) $(ALL_LIBS) -lopp_hip -lz
+		$(ALL_INC) $(HIP_INC) $(ALL_LIBS) -lopp_hip -lz $(HIP_LIB)
 
 # ------------------------------------------------------------------------------------------
 advec_opp+hip_mpi:
-	 $(HIPCC) $(HIPCCFLAGS) -DUSE_HIP -DUSE_MPI -c advec_opp.cpp -o $(OBJ)/advec_opp_hip_mpi.o $(ALL_INC) $(HIP_INC)
+	 $(MPICPP) $(CPPFLAGS) -DUSE_HIP -DUSE_MPI -c advec_opp.cpp -o $(OBJ)/advec_opp_hip_mpi.o $(ALL_INC) $(HIP_INC)
 opp_kernels+hip_mpi:
 	 $(HIPCC) $(HIPCCFLAGS) -DUSE_HIP -DUSE_MPI -c hip/opp_kernels.cpp -o $(OBJ)/opp_kernels_hip_mpi.o $(ALL_INC) $(HIP_INC)
 hip_mpi: mklib advec_opp+hip_mpi opp_kernels+hip_mpi
-	 $(HIPCC) $(HIPCCFLAGS) -o $(BIN)/hip_mpi \
+	 $(MPICPP) $(CPPFLAGS) -o $(BIN)/hip_mpi \
 		$(OBJ)/advec_opp_hip_mpi.o \
 		$(OBJ)/opp_kernels_hip_mpi.o \
-		$(ALL_INC) $(HIP_INC) $(ALL_LIBS) -lopp_hip_mpi -lmpi -lz
+		$(ALL_INC) $(HIP_INC) $(ALL_LIBS) -lopp_hip_mpi -lmpi -lz $(HIP_LIB)
 
 # ------------------------------------------------------------------------------------------
 clean: 

app_neso_advection_cg/sycl/move_kernel_loop.hpp

@@ -193,7 +193,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 // c_pos_ll | OPP_READ

opp_lib/Makefile

@@ -85,7 +85,7 @@ CUDA_INC   += -I$(CUDA_INSTALL_PATH)/include
 # for demos with gcc
 # CUDA_INC   = -I$(CUDA_INSTALL_PATH)/include -I$(MPI_INSTALL_PATH)/include
 
-HIP_INC   += -I$(ROCM_INSTALL_DIR)/include -I$(ROCM_THRUST_DIR)/include -I$(ROCM_PRIM_DIR)/include
+HIP_INC   += -I$(ROCM_INSTALL_DIR)/include -I$(ROCM_THRUST_DIR)/include -I$(ROCM_PRIM_DIR)/include -I$(MPI_INSTALL_PATH)/include
 
 .PHONY: clean mklib
 
@@ -154,7 +154,8 @@ SRC_SYCL_CPP = 	$(SRC)/sycl/opp_sycl.cpp \
 				$(SRC)/sycl/opp_sycl_halo.cpp \
 				$(SRC)/core/opp_lib_core.cpp \
 				$(SRC)/core/opp_util.cpp \
-				$(SRC)/core/opp_params.cpp
+				$(SRC)/core/opp_params.cpp \
+				$(SRC)/sycl/opp_direct_hop_sycl.cpp
 OBJ_SYCL_CPP =  $(patsubst $(SRC)/core/%.cpp,$(OBJ)/%+sycl.o,$(SRC_SYCL_CPP)) \
 				$(patsubst $(SRC)/sycl/%.cpp,$(OBJ)/%+sycl.o,$(SRC_SYCL_CPP))
 $(OBJ)/%+sycl.o: $(SRC)/core/%.cpp
@@ -173,7 +174,8 @@ SRC_HIP_CPP = 	$(SRC)/core/opp_lib_core.cpp \
 				$(SRC)/core/opp_params.cpp \
 				$(SRC)/hip/opp_hip.cpp \
 				$(SRC)/hip/opp_particle_sorter.cpp \
-				$(SRC)/hip/opp_hip_halo.cpp 
+				$(SRC)/hip/opp_hip_halo.cpp \
+				$(SRC)/hip/opp_direct_hop_hip.cpp
 OBJ_HIP_CPP =  	$(patsubst $(SRC)/core/%.cpp,$(OBJ)/%+hip.o,$(SRC_HIP_CPP)) \
 				$(patsubst $(SRC)/hip/%.cpp,$(OBJ)/%+hip.o,$(SRC_HIP_CPP))
 $(OBJ)/%+hip.o: $(SRC)/core/%.cpp
@@ -223,7 +225,8 @@ SRC_SYCL_MPI_CPP = 	$(SRC)/sycl/opp_sycl.cpp \
 					$(SRC)/mpi/opp_mpi_core.cpp \
 					$(SRC)/mpi/opp_mpi_particle_comm.cpp \
 					$(SRC)/mpi/opp_mpi_double_ind_reducs.cpp \
-					$(SRC)/mpi/opp_mpi_utils.cpp
+					$(SRC)/mpi/opp_mpi_utils.cpp \
+					$(SRC)/sycl/opp_direct_hop_sycl.cpp
 OBJ_SYCL_MPI_CPP =  $(patsubst $(SRC)/core/%.cpp,$(OBJ)/%+sycl_mpi.o,$(SRC_SYCL_MPI_CPP)) \
 					$(patsubst $(SRC)/sycl/%.cpp,$(OBJ)/%+sycl_mpi.o,$(SRC_SYCL_MPI_CPP)) \
 					$(patsubst $(SRC)/mpi/%.cpp,$(OBJ)/%+sycl_mpi.o,$(SRC_SYCL_MPI_CPP))
@@ -280,7 +283,8 @@ SRC_HIP_MPI_CPP = 	$(SRC)/core/opp_lib_core.cpp \
 					$(SRC)/mpi/opp_mpi_utils.cpp \
 					$(SRC)/hip/opp_hip.cpp \
 					$(SRC)/hip/opp_hip_halo.cpp \
-					$(SRC)/hip/opp_particle_sorter.cpp
+					$(SRC)/hip/opp_particle_sorter.cpp \
+					$(SRC)/hip/opp_direct_hop_hip.cpp
 OBJ_HIP_MPI_CPP = 	$(patsubst $(SRC)/core/%.cpp,$(OBJ)/%+hip_mpi.o,$(SRC_HIP_MPI_CPP)) \
 					$(patsubst $(SRC)/mpi/%.cpp,$(OBJ)/%+hip_mpi.o,$(SRC_HIP_MPI_CPP)) \
 					$(patsubst $(SRC)/hip/%.cpp,$(OBJ)/%+hip_mpi.o,$(SRC_HIP_MPI_CPP))

opp_lib/include/opp_hip.h

@@ -53,6 +53,14 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
     #include <opp_mpi_core.h>
 #endif
 
+struct opp_dh_indices {
+    OPP_INT* move_count = nullptr;
+    OPP_INT* part_indices = nullptr;
+    OPP_INT* rank_indices = nullptr;
+    OPP_INT* cell_indices = nullptr;
+    OPP_INT capacity = 0;
+};
+
 #define cutilSafeCall(err) __hipSafeCall(err, __FILE__, __LINE__)
 #define cutilCheckMsg(msg) __cutilCheckMsg(msg, __FILE__, __LINE__)
 
@@ -92,11 +100,19 @@ extern std::map<int, thrust::device_vector<char>> send_data;
 extern std::map<int, thrust::device_vector<char>> recv_data;
 
 // arrays for global constants and reductions
-extern int OPP_consts_bytes;
-extern int OPP_reduct_bytes;
+extern OPP_INT OPP_consts_bytes;
+extern OPP_INT OPP_reduct_bytes;
 extern char *OPP_reduct_h, *OPP_reduct_d;
 extern char *OPP_consts_h, *OPP_consts_d;
 
+extern char opp_move_status_flag;
+extern bool opp_move_hop_iter_one_flag;
+extern OPP_INT* opp_p2c;
+extern OPP_INT* opp_c2c;
+
+extern opp_dh_indices dh_indices_d;
+extern opp_dh_indices dh_indices_h;
+
 //*************************************************************************************************
 
 void __hipSafeCall(hipError_t err, const char *file, const int line);

opp_lib/include/opp_sycl.h

@@ -45,6 +45,14 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
     #include <opp_mpi_core.h>
 #endif
 
+struct opp_dh_indices {
+    OPP_INT* move_count = nullptr;
+    OPP_INT* part_indices = nullptr;
+    OPP_INT* rank_indices = nullptr;
+    OPP_INT* cell_indices = nullptr;
+    OPP_INT capacity = 0;
+};
+
 constexpr bool debug_mem = false;
 constexpr bool debugger = false;
 constexpr int opp_const_threads_per_block = 192;
@@ -90,6 +98,14 @@ extern char *OPP_consts_h, *OPP_consts_d;
 
 extern std::vector<char*> opp_consts;
 
+extern char opp_move_status_flag;
+extern bool opp_move_hop_iter_one_flag;
+extern OPP_INT* opp_p2c;
+extern OPP_INT* opp_c2c;
+
+extern opp_dh_indices dh_indices_d;
+extern opp_dh_indices dh_indices_h;
+
 //*************************************************************************************************
 void opp_sycl_init(int argc, char **argv);
 void opp_sycl_exit();

opp_lib/src/cuda/opp_cuda_utils.cu

@@ -36,7 +36,7 @@ void opp_create_dat_device_arrays(opp_dat dat, bool create_new)
 {
     if (OPP_DBG) opp_printf("opp_create_dat_device_arrays", "%s %s", dat->name, dat->type);
 
-    char* temp_char_d = nullptr;
+    // char* temp_char_d = nullptr;
     const size_t capacity = (size_t)dat->set->set_capacity * dat->dim;
 
     if (strcmp(dat->type, "double") == 0) {
@@ -50,7 +50,7 @@ void opp_create_dat_device_arrays(opp_dat dat, bool create_new)
             dat->data_d = (char*)opp_get_dev_raw_ptr(*(dat->thrust_real));
 
             dat->thrust_real_sort = new thrust::device_vector<OPP_REAL>(capacity);
-            temp_char_d = (char*)opp_get_dev_raw_ptr(*(dat->thrust_real_sort));
+            dat->data_swap_d = (char*)opp_get_dev_raw_ptr(*(dat->thrust_real_sort));
         } 
     } 
     else if (strcmp(dat->type, "int") == 0 ) {
@@ -64,7 +64,7 @@ void opp_create_dat_device_arrays(opp_dat dat, bool create_new)
             dat->data_d = (char*)opp_get_dev_raw_ptr(*(dat->thrust_int));
 
             dat->thrust_int_sort = new thrust::device_vector<OPP_INT>(capacity);
-            temp_char_d = (char*)opp_get_dev_raw_ptr(*(dat->thrust_int_sort));
+            dat->data_swap_d = (char*)opp_get_dev_raw_ptr(*(dat->thrust_int_sort));
         } 
     }
     else {
@@ -75,7 +75,7 @@ void opp_create_dat_device_arrays(opp_dat dat, bool create_new)
 
     if (OPP_DBG) 
         opp_printf("opp_create_dat_device_arrays", "Device array of dat [%s][%p][%p] Capacity [%d]", 
-                        dat->name, dat->data_d, temp_char_d, capacity);
+                        dat->name, dat->data_d, dat->data_swap_d, capacity);
 }
 
 //****************************************

opp_lib/src/cuda/opp_particle_sorter.cu

@@ -141,6 +141,7 @@ void sort_dat_according_to_index_int(opp_dat dat, const thrust::device_vector<in
         dat->thrust_int_sort = tmp;
 
         dat->data_d = (char*)opp_get_dev_raw_ptr(*(dat->thrust_int));
+        dat->data_swap_d = (char*)opp_get_dev_raw_ptr(*(dat->thrust_real_int));
     }
 }
 
@@ -178,6 +179,7 @@ void sort_dat_according_to_index_double(opp_dat dat, const thrust::device_vector
         dat->thrust_real_sort = tmp;
 
         dat->data_d = (char*)opp_get_dev_raw_ptr(*(dat->thrust_real));
+        dat->data_swap_d = (char*)opp_get_dev_raw_ptr(*(dat->thrust_real_sort));
     }
 }