Merge branch 'stp/cmake_refactor_to_merge' of github.com:ExCALIBUR-NE…

…PTUNE/NESO into stp/cmake_refactor_to_merge

include/mesh.hpp

@@ -28,17 +28,13 @@ class Mesh {
   int nintervals;
   // number of grid points (including periodic point)
   int nmesh;
-  sycl::buffer<int, 1> nmesh_d;
   // grid spacing
   double dx;
-  sycl::buffer<double, 1> dx_d;
 
   // box length in units of Debye length
   double normalized_box_length;
   // mesh point vector
   std::vector<double> mesh;
-  // mesh points (device buffer)
-  sycl::buffer<double, 1> mesh_d;
   // mesh point vector staggered at half points
   std::vector<double> mesh_staggered;
   // Fourier wavenumbers corresponding to mesh
@@ -47,15 +43,11 @@ class Mesh {
   std::vector<double> poisson_factor;
   // Factor to use in combined field solve and E = -Grad(phi)
   std::vector<Complex> poisson_E_factor;
-  sycl::buffer<Complex, 1> poisson_E_factor_d;
 
   // charge density
   std::vector<double> charge_density;
-  sycl::buffer<double, 1> charge_density_d;
   // electric field
   std::vector<double> electric_field;
-  // electric field (device)
-  sycl::buffer<double, 1> electric_field_d;
   // electric field on a staggered grid
   std::vector<double> electric_field_staggered;
   // electrostatic potential
@@ -64,33 +56,6 @@ class Mesh {
   // Calculate a particle's contribution to the electric field
   double evaluate_electric_field(const double x);
 
-  /*
-   * Evaluate the electric field at x grid points by
-   * interpolating onto the grid
-   * SYCL note: this is a copy of evaluate_electric_field, but able to be called
-   * in sycl. This should become evaluate_electric_field eventually
-   */
-  inline double
-  sycl_evaluate_electric_field(sycl::accessor<double> mesh_d,
-                               sycl::accessor<double> electric_field_d,
-                               double x) {
-
-    // Find grid cell that x is in
-    int index = sycl_get_left_index(x, mesh_d);
-
-    // now x is in the cell ( mesh[index-1], mesh[index] )
-
-    double cell_width = mesh_d[index + 1] - mesh_d[index];
-    double distance_into_cell = x - mesh_d[index];
-
-    // r is the proportion if the distance into the cell that the particle is at
-    // e.g. midpoint => r = 0.5
-    double r = distance_into_cell / cell_width;
-
-    return (1.0 - r) * electric_field_d[index] +
-           r * electric_field_d[index + 1];
-  };
-
   // Deposit particle onto mesh
   void deposit(Plasma &plasma);
   void sycl_deposit(sycl::queue &Q, Plasma &plasma);
@@ -118,16 +83,46 @@ class Mesh {
   // Given a point x and a grid, find the indices of the grid points
   // either side of x
   int get_left_index(const double x, const std::vector<double> mesh);
+};
 
-  inline int sycl_get_left_index(const double x,
-                                 const sycl::accessor<double> mesh_d) {
+namespace Mesh1D {
 
-    int index = 0;
-    while (mesh_d[index + 1] < x and index < int(mesh.size())) {
-      index++;
-    };
-    return index;
-  }
-};
+template <typename T>
+inline int sycl_get_left_index(const double x, const T &mesh_d,
+                               const int mesh_size) {
+
+  int index = 0;
+  while ((mesh_d[index + 1] < x) and (index < mesh_size)) {
+    index++;
+  };
+  return index;
+}
+
+/*
+ * Evaluate the electric field at x grid points by
+ * interpolating onto the grid
+ * SYCL note: this is a copy of evaluate_electric_field, but able to be called
+ * in sycl. This should become evaluate_electric_field eventually
+ */
+inline double sycl_evaluate_electric_field(
+    const sycl::accessor<double> &mesh_d, const int mesh_size,
+    const sycl::accessor<double> &electric_field_d, double x) {
+
+  // Find grid cell that x is in
+  int index = sycl_get_left_index(x, mesh_d, mesh_size);
+
+  // now x is in the cell ( mesh[index-1], mesh[index] )
+
+  double cell_width = mesh_d[index + 1] - mesh_d[index];
+  double distance_into_cell = x - mesh_d[index];
+
+  // r is the proportion if the distance into the cell that the particle is at
+  // e.g. midpoint => r = 0.5
+  double r = distance_into_cell / cell_width;
+
+  return (1.0 - r) * electric_field_d[index] + r * electric_field_d[index + 1];
+}
+
+} // namespace Mesh1D
 
 #endif // __MESH_H__

include/nektar_interface/function_basis_evaluation.hpp

@@ -99,8 +99,7 @@ class FunctionEvaluateBasis : public BasisEvaluateBase<T> {
     const int k_ndim = evaluation_type.get_ndim();
 
     sycl::range<2> cell_iterset_range{static_cast<size_t>(cells_iterset_size),
-                                      static_cast<size_t>(outer_size) *
-                                          static_cast<size_t>(local_size)};
+                                      static_cast<size_t>(outer_size)};
     sycl::range<2> local_iterset{1, local_size};
 
     auto event_loop = this->sycl_target->queue.submit([&](sycl::handler &cgh) {

include/nektar_interface/function_basis_projection.hpp

@@ -99,8 +99,7 @@ template <typename T> class FunctionProjectBasis : public BasisEvaluateBase<T> {
     const int k_ndim = project_type.get_ndim();
 
     sycl::range<2> cell_iterset_range{static_cast<size_t>(cells_iterset_size),
-                                      static_cast<size_t>(outer_size) *
-                                          static_cast<size_t>(local_size)};
+                                      static_cast<size_t>(outer_size)};
     sycl::range<2> local_iterset{1, local_size};
 
     auto event_loop = this->sycl_target->queue.submit([&](sycl::handler &cgh) {

include/species.hpp

@@ -30,19 +30,10 @@ class Species {
   double vth;
   // particle position array
   std::vector<double> x;
-  // particle positions (device)
-  sycl::buffer<double, 1> x_d;
   // particle velocity structure of arrays
   Velocity v;
-  // particle x velocity (device)
-  sycl::buffer<double, 1> vx_d;
-  // particle y velocity (device)
-  sycl::buffer<double, 1> vy_d;
-  // particle z velocity (device)
-  sycl::buffer<double, 1> vz_d;
   // charge density of species (if adiabatic)
   double charge_density;
-  sycl::buffer<double, 1> charge_density_d;
   // particle position array at
   // next timestep
   std::vector<double> xnew;
@@ -51,8 +42,6 @@ class Species {
   std::vector<double> vnew;
   // particle weight
   std::vector<double> w;
-  // particle weights (device)
-  sycl::buffer<double, 1> w_d;
   // particle pusher
   void push(sycl::queue &q, Mesh *mesh);
   // set array sizes for particle properties
@@ -62,8 +51,6 @@ class Species {
   // Coefficients for particle pusher
   double dx_coef;
   double dv_coef;
-  sycl::buffer<double, 1> dx_coef_d;
-  sycl::buffer<double, 1> dv_coef_d;
 };
 
 #endif // __SPECIES_H__

solvers/Electrostatic2D3V/EquationSystems/PoissonPIC.cpp

@@ -1,4 +1,4 @@
-#include "PoissonPIC.h"
+#include "PoissonPIC.hpp"
 
 using namespace std;
 

solvers/Electrostatic2D3V/EquationSystems/PoissonPIC.hpp

@@ -0,0 +1,53 @@
+#ifndef NEKTAR_SOLVERS_EQUATIONSYSTEMS_POISSON_PIC_H
+#define NEKTAR_SOLVERS_EQUATIONSYSTEMS_POISSON_PIC_H
+
+#include <map>
+#include <string>
+
+#include <SolverUtils/EquationSystem.h>
+using namespace Nektar::SolverUtils;
+
+namespace Nektar {
+class PoissonPIC : public EquationSystem {
+public:
+  std::map<std::string, int> field_to_index;
+
+  friend class MemoryManager<PoissonPIC>;
+
+  /// Creates an instance of this class
+  static EquationSystemSharedPtr
+  create(const LibUtilities::SessionReaderSharedPtr &pSession,
+         const SpatialDomains::MeshGraphSharedPtr &pGraph) {
+    EquationSystemSharedPtr p =
+        MemoryManager<PoissonPIC>::AllocateSharedPtr(pSession, pGraph);
+    p->InitObject();
+    return p;
+  }
+  /// Name of class
+  static std::string className1;
+  static std::string className2;
+
+  virtual ~PoissonPIC();
+  /**
+   *  Helper function to map from field name to field indices.
+   *
+   *  @param name Field name (probably either "u" or "rho").
+   *  @returns index (probably 0 or 1).
+   */
+  int GetFieldIndex(const std::string name);
+
+protected:
+  StdRegions::ConstFactorMap m_factors;
+  PoissonPIC(const LibUtilities::SessionReaderSharedPtr &pSession,
+             const SpatialDomains::MeshGraphSharedPtr &pGraph);
+
+  virtual void v_InitObject(bool DeclareFields = true);
+  virtual void v_DoSolve();
+  virtual void v_GenerateSummary(SolverUtils::SummaryList &s);
+
+private:
+  virtual Array<OneD, bool> v_GetSystemSingularChecks();
+};
+} // namespace Nektar
+
+#endif

solvers/Electrostatic2D3V/ParticleSystems/poisson_particle_coupling.hpp

@@ -21,7 +21,7 @@
 #include <random>
 #include <string>
 
-#include "../EquationSystems/PoissonPIC.h"
+#include "../EquationSystems/PoissonPIC.hpp"
 #include "charged_particles.hpp"
 
 using namespace Nektar;

solvers/SimpleSOL/Diagnostics/mass_conservation.hpp

@@ -1,23 +1,21 @@
 #ifndef __SIMPLESOL_MASS_CONSERVATION_H_
 #define __SIMPLESOL_MASS_CONSERVATION_H_
 
+#include <LibUtilities/BasicUtils/ErrorUtil.hpp>
 #include <memory>
 #include <mpi.h>
 #include <neso_particles.hpp>
-using namespace NESO;
-using namespace NESO::Particles;
-
-#include <LibUtilities/BasicUtils/ErrorUtil.hpp>
-using namespace Nektar;
 
 #include "../ParticleSystems/neutral_particles.hpp"
 
 #include <fstream>
 #include <iostream>
 
+namespace LU = Nektar::LibUtilities;
+
 template <typename T> class MassRecording {
 protected:
-  const LibUtilities::SessionReaderSharedPtr session;
+  const LU::SessionReaderSharedPtr session;
   std::shared_ptr<NeutralParticleSystem> particle_sys;
   std::shared_ptr<T> rho;
 
@@ -30,7 +28,7 @@ template <typename T> class MassRecording {
   std::ofstream fh;
 
 public:
-  MassRecording(const LibUtilities::SessionReaderSharedPtr session,
+  MassRecording(const LU::SessionReaderSharedPtr session,
                 std::shared_ptr<NeutralParticleSystem> particle_sys,
                 std::shared_ptr<T> rho)
       : session(session), particle_sys(particle_sys), rho(rho),