Merge pull request #14 from OP-DSL/opp_cg

Merge opp_cg to main
OP-DSL · Jun 15, 2024 · 1dfa687 · 1dfa687
2 parents ed53800 + 1fc85fb
commit 1dfa687
Show file tree

Hide file tree

Showing 90 changed files with 1,310 additions and 804 deletions.
diff --git a/README.md b/README.md
@@ -8,7 +8,10 @@ This repository contains the implementation of the code translation tools and ru
  * `app_<app_name>_cg`: Copy of the example applications that are already being code generated. 
  * `scripts/source`: Some example source files used during implementation
  * `scripts/batch`: Some example slurm batch files used during implementation
- * `handcoded`: Some hand-coded application code, written prior code generation
+ * `app_handcoded`: Some hand-coded application code, written prior code generation
+
+## Documentation
+Documentation is available on [Read the Docs](https://zamanlantra.github.io/git_docs_test/index.html).
 
 ## Dependencies
 OP-PIC has a variety of toolchain dependencies that you will likely be able to obtain from your package manager or programming environment:
@@ -39,7 +42,7 @@ In addition, there are a few optional library dependencies that you will likely
  * Compile the required application version using the make command. (`make` followed by the required parallelization). 
 
  For example, `cd app_cabanapic; python3 $OPP_TRANSLATOR -v -I$OPP_PATH/include/ --file_paths cabanapic.cpp; make cuda_mpi`.
- A detail explanation can be found in the readme file of opp_translator folder and the required application folder.
+ A detail explanation can be found in the readme file of `opp_translator` folder and the required application folder.
 
  In addition, `app_<app_name>_cg` will additionally include code-generated files, ready to compile directly using make commands.
 

diff --git a/app_cabanapic/Makefile b/app_cabanapic/Makefile
@@ -92,6 +92,14 @@ mpi: mklib
 			$(OBJ)/opp_kernels_mpi.o \
 			$(ALL_INC) $(ALL_LIBS) -lopp_mpi
 
+omp_mpi: mklib
+	     $(MPICPP) $(CPPFLAGS) -DUSE_MPI -fopenmp -DUSE_OMP -c cabana_opp.cpp -o $(OBJ)/cabana_opp_omp_mpi.o $(ALL_INC)
+	     $(MPICPP) $(CPPFLAGS) -DUSE_MPI -fopenmp -DUSE_OMP -c omp/opp_kernels.cpp -o $(OBJ)/opp_kernels_omp_mpi.o $(ALL_INC)
+	     $(MPICPP) $(CPPFLAGS) -DUSE_MPI -fopenmp -o $(BIN)/omp_mpi \
+			$(OBJ)/cabana_opp_omp_mpi.o \
+			$(OBJ)/opp_kernels_omp_mpi.o \
+			$(ALL_INC) $(ALL_LIBS) -lopp_omp_mpi
+
 mpi_hdf5: mklib
 		 $(MPICPP) $(CPPFLAGS) -DUSE_MPI -c cabana_op_hdf5.cpp -o $(OBJ)/cabana_op_hdf5.o $(ALL_INC)
 		 $(MPICPP) $(CPPFLAGS) -DUSE_MPI -c mpi/mpikernels.cpp -o $(OBJ)/mpikernels.o $(ALL_INC)

diff --git a/app_cabanapic_cg/Makefile b/app_cabanapic_cg/Makefile
@@ -92,6 +92,14 @@ mpi: mklib
 			$(OBJ)/opp_kernels_mpi.o \
 			$(ALL_INC) $(ALL_LIBS) -lopp_mpi
 
+omp_mpi: mklib
+	     $(MPICPP) $(CPPFLAGS) -DUSE_MPI -fopenmp -DUSE_OMP -c cabana_opp.cpp -o $(OBJ)/cabana_opp_omp_mpi.o $(ALL_INC)
+	     $(MPICPP) $(CPPFLAGS) -DUSE_MPI -fopenmp -DUSE_OMP -c omp/opp_kernels.cpp -o $(OBJ)/opp_kernels_omp_mpi.o $(ALL_INC)
+	     $(MPICPP) $(CPPFLAGS) -DUSE_MPI -fopenmp -o $(BIN)/omp_mpi \
+			$(OBJ)/cabana_opp_omp_mpi.o \
+			$(OBJ)/opp_kernels_omp_mpi.o \
+			$(ALL_INC) $(ALL_LIBS) -lopp_omp_mpi
+
 mpi_hdf5: mklib
 		 $(MPICPP) $(CPPFLAGS) -DUSE_MPI -c cabana_op_hdf5.cpp -o $(OBJ)/cabana_op_hdf5.o $(ALL_INC)
 		 $(MPICPP) $(CPPFLAGS) -DUSE_MPI -c mpi/mpikernels.cpp -o $(OBJ)/mpikernels.o $(ALL_INC)

diff --git a/app_cabanapic_cg/cuda/interpolate_mesh_fields_kernel_loop.hpp b/app_cabanapic_cg/cuda/interpolate_mesh_fields_kernel_loop.hpp
@@ -18,6 +18,12 @@ __constant__ OPP_INT opp_k1_map0_stride_d;
 
 
 namespace opp_k1 {
+enum Dim {
+    x = 0,
+    y = 1,
+    z = 2,
+};
+
 enum CellInterp {
     ex = 0,
     dexdy,
@@ -39,12 +45,6 @@ enum CellInterp {
     dcbzdz,
 };
 
-enum Dim {
-    x = 0,
-    y = 1,
-    z = 2,
-};
-
 __device__ inline void interpolate_mesh_fields_kernel(
     const double* cell0_e,
     const double* cell0_b,

diff --git a/app_cabanapic_cg/cuda/move_deposit_kernel_loop.hpp b/app_cabanapic_cg/cuda/move_deposit_kernel_loop.hpp
@@ -26,27 +26,6 @@ enum CellAcc {
     jfz = 2 * 4,
 };
 
-enum CellInterp {
-    ex = 0,
-    dexdy,
-    dexdz,
-    d2exdydz,
-    ey,
-    deydz,
-    deydx,
-    d2eydzdx,
-    ez,
-    dezdx,
-    dezdy,
-    d2ezdxdy,
-    cbx,
-    dcbxdx,
-    cby,
-    dcbydy,
-    cbz,
-    dcbzdz,
-};
-
 __device__ inline void weight_current_to_accumulator_kernel(
         double* cell_acc,
         const double* q,
@@ -82,6 +61,27 @@ enum Dim {
     z = 2,
 };
 
+enum CellInterp {
+    ex = 0,
+    dexdy,
+    dexdz,
+    d2exdydz,
+    ey,
+    deydz,
+    deydx,
+    d2eydzdx,
+    ez,
+    dezdx,
+    dezdy,
+    d2ezdxdy,
+    cbx,
+    dcbxdx,
+    cby,
+    dcbydy,
+    cbz,
+    dcbzdz,
+};
+
 __device__ inline void move_deposit_kernel(
     char& opp_move_status_flag, const bool opp_move_hop_iter_one_flag, // Added by code-gen
     const OPP_INT* opp_c2c, OPP_INT* opp_p2c, // Added by code-gen

diff --git a/app_cabanapic_cg/hip/interpolate_mesh_fields_kernel_loop.hpp b/app_cabanapic_cg/hip/interpolate_mesh_fields_kernel_loop.hpp
@@ -18,6 +18,12 @@ __constant__ OPP_INT opp_k1_map0_stride_d;
 
 
 namespace opp_k1 {
+enum Dim {
+    x = 0,
+    y = 1,
+    z = 2,
+};
+
 enum CellInterp {
     ex = 0,
     dexdy,
@@ -39,12 +45,6 @@ enum CellInterp {
     dcbzdz,
 };
 
-enum Dim {
-    x = 0,
-    y = 1,
-    z = 2,
-};
-
 __device__ inline void interpolate_mesh_fields_kernel(
     const double* cell0_e,
     const double* cell0_b,

diff --git a/app_cabanapic_cg/hip/move_deposit_kernel_loop.hpp b/app_cabanapic_cg/hip/move_deposit_kernel_loop.hpp
@@ -26,27 +26,6 @@ enum CellAcc {
     jfz = 2 * 4,
 };
 
-enum CellInterp {
-    ex = 0,
-    dexdy,
-    dexdz,
-    d2exdydz,
-    ey,
-    deydz,
-    deydx,
-    d2eydzdx,
-    ez,
-    dezdx,
-    dezdy,
-    d2ezdxdy,
-    cbx,
-    dcbxdx,
-    cby,
-    dcbydy,
-    cbz,
-    dcbzdz,
-};
-
 __device__ inline void weight_current_to_accumulator_kernel(
         double* cell_acc,
         const double* q,
@@ -82,6 +61,27 @@ enum Dim {
     z = 2,
 };
 
+enum CellInterp {
+    ex = 0,
+    dexdy,
+    dexdz,
+    d2exdydz,
+    ey,
+    deydz,
+    deydx,
+    d2eydzdx,
+    ez,
+    dezdx,
+    dezdy,
+    d2ezdxdy,
+    cbx,
+    dcbxdx,
+    cby,
+    dcbydy,
+    cbz,
+    dcbzdz,
+};
+
 __device__ inline void move_deposit_kernel(
     char& opp_move_status_flag, const bool opp_move_hop_iter_one_flag, // Added by code-gen
     const OPP_INT* opp_c2c, OPP_INT* opp_p2c, // Added by code-gen

diff --git a/app_cabanapic_cg/mpi/interpolate_mesh_fields_kernel_loop.hpp b/app_cabanapic_cg/mpi/interpolate_mesh_fields_kernel_loop.hpp
@@ -4,6 +4,12 @@
 //*********************************************
 
 namespace opp_k1 {
+enum Dim {
+    x = 0,
+    y = 1,
+    z = 2,
+};
+
 enum CellInterp {
     ex = 0,
     dexdy,
@@ -25,12 +31,6 @@ enum CellInterp {
     dcbzdz,
 };
 
-enum Dim {
-    x = 0,
-    y = 1,
-    z = 2,
-};
-
 inline void interpolate_mesh_fields_kernel(
     const double* cell0_e,
     const double* cell0_b,

diff --git a/app_cabanapic_cg/mpi/move_deposit_kernel_loop.hpp b/app_cabanapic_cg/mpi/move_deposit_kernel_loop.hpp
@@ -10,27 +10,6 @@ enum CellAcc {
     jfz = 2 * 4,
 };
 
-enum CellInterp {
-    ex = 0,
-    dexdy,
-    dexdz,
-    d2exdydz,
-    ey,
-    deydz,
-    deydx,
-    d2eydzdx,
-    ez,
-    dezdx,
-    dezdy,
-    d2ezdxdy,
-    cbx,
-    dcbxdx,
-    cby,
-    dcbydy,
-    cbz,
-    dcbzdz,
-};
-
 inline void weight_current_to_accumulator_kernel(
         double* cell_acc,
         const double* q,
@@ -66,6 +45,27 @@ enum Dim {
     z = 2,
 };
 
+enum CellInterp {
+    ex = 0,
+    dexdy,
+    dexdz,
+    d2exdydz,
+    ey,
+    deydz,
+    deydx,
+    d2eydzdx,
+    ez,
+    dezdx,
+    dezdy,
+    d2ezdxdy,
+    cbx,
+    dcbxdx,
+    cby,
+    dcbydy,
+    cbz,
+    dcbzdz,
+};
+
 inline void move_deposit_kernel(
     double* part_vel,
     double* part_pos,
@@ -341,7 +341,7 @@ void opp_particle_move__move_deposit_kernel(opp_set set, opp_map c2c_map, opp_ma
     };
 
     // ----------------------------------------------------------------------------
-    opp_init_particle_move(set, 0, nullptr);
+    opp_init_particle_move(set, nargs, args);
 
 
     opp_profiler->start("Mv_AllMv0");
@@ -369,7 +369,7 @@ void opp_particle_move__move_deposit_kernel(opp_set set, opp_map c2c_map, opp_ma
         const std::string profName = std::string("Mv_AllMv") + std::to_string(OPP_comm_iteration);
         opp_profiler->start(profName);
 
-        opp_init_particle_move(set, 0, nullptr);
+        opp_init_particle_move(set, nargs, args);
 
         // check whether particle is within cell, and if not move between cells within the MPI rank, mark for neighbour comm
         opp_profiler->start("move_deposit_kernel_only");

diff --git a/app_cabanapic_cg/omp/accumulate_current_to_cells_kernel_loop.hpp b/app_cabanapic_cg/omp/accumulate_current_to_cells_kernel_loop.hpp
@@ -88,8 +88,8 @@ void opp_par_loop_all__accumulate_current_to_cells_kernel(opp_set set, opp_itera
     #pragma omp parallel for 
     for (int thr = 0; thr < nthreads; thr++)
     {
-        const size_t start  = (iter_size * thr) / nthreads;
-        const size_t finish = (iter_size * (thr+1)) / nthreads;
+        const size_t start  = ((size_t)iter_size * thr) / nthreads;
+        const size_t finish = ((size_t)iter_size * (thr+1)) / nthreads;
 
         for (size_t n = start; n < finish; n++)
         { 

diff --git a/app_cabanapic_cg/omp/advance_e_kernel_loop.hpp b/app_cabanapic_cg/omp/advance_e_kernel_loop.hpp
@@ -73,8 +73,8 @@ void opp_par_loop_all__advance_e_kernel(opp_set set, opp_iterate_type,
     #pragma omp parallel for 
     for (int thr = 0; thr < nthreads; thr++)
     {
-        const size_t start  = (iter_size * thr) / nthreads;
-        const size_t finish = (iter_size * (thr+1)) / nthreads;
+        const size_t start  = ((size_t)iter_size * thr) / nthreads;
+        const size_t finish = ((size_t)iter_size * (thr+1)) / nthreads;
 
         for (size_t n = start; n < finish; n++)
         { 

diff --git a/app_cabanapic_cg/omp/compute_energy_kernel_loop.hpp b/app_cabanapic_cg/omp/compute_energy_kernel_loop.hpp
@@ -59,8 +59,8 @@ void opp_par_loop_all__compute_energy_kernel(opp_set set, opp_iterate_type,
     #pragma omp parallel for 
     for (int thr = 0; thr < nthreads; thr++)
     {
-        const size_t start  = (iter_size * thr) / nthreads;
-        const size_t finish = (iter_size * (thr+1)) / nthreads;
+        const size_t start  = ((size_t)iter_size * thr) / nthreads;
+        const size_t finish = ((size_t)iter_size * (thr+1)) / nthreads;
 
         for (size_t n = start; n < finish; n++)
         { 

diff --git a/app_cabanapic_cg/omp/get_max_x_values_kernel_loop.hpp b/app_cabanapic_cg/omp/get_max_x_values_kernel_loop.hpp
@@ -73,8 +73,8 @@ void opp_par_loop_all__get_max_x_values_kernel(opp_set set, opp_iterate_type,
     #pragma omp parallel for 
     for (int thr = 0; thr < nthreads; thr++)
     {
-        const size_t start  = (iter_size * thr) / nthreads;
-        const size_t finish = (iter_size * (thr+1)) / nthreads;
+        const size_t start  = ((size_t)iter_size * thr) / nthreads;
+        const size_t finish = ((size_t)iter_size * (thr+1)) / nthreads;
 
         for (size_t n = start; n < finish; n++)
         { 

diff --git a/app_cabanapic_cg/omp/half_advance_b_kernel_loop.hpp b/app_cabanapic_cg/omp/half_advance_b_kernel_loop.hpp
@@ -67,8 +67,8 @@ void opp_par_loop_all__half_advance_b_kernel(opp_set set, opp_iterate_type,
     #pragma omp parallel for 
     for (int thr = 0; thr < nthreads; thr++)
     {
-        const size_t start  = (iter_size * thr) / nthreads;
-        const size_t finish = (iter_size * (thr+1)) / nthreads;
+        const size_t start  = ((size_t)iter_size * thr) / nthreads;
+        const size_t finish = ((size_t)iter_size * (thr+1)) / nthreads;
 
         for (size_t n = start; n < finish; n++)
         {