urgent fix to gcc related build errors

include/neso_particles/containers/local_memory_block.hpp

@@ -25,9 +25,7 @@ class LocalMemoryBlock {
    *
    * @param size Number of bytes required per particle.
    */
-  LocalMemoryBlock(const std::size_t size) : size(size) {
-    NESOASSERT(size >= 0, "Invalid local size passed.");
-  }
+  LocalMemoryBlock(const std::size_t size) : size(size) {}
 
   /**
    * Set the local memory required to a new number of bytes.

include/neso_particles/device_limits.hpp

@@ -6,51 +6,63 @@
 
 namespace NESO::Particles {
 
-class DeviceLimits {
-protected:
+namespace Private {
+
+struct WorkGroupLimits {
   sycl::range<1> max_global_workgroup_1;
   sycl::range<2> max_global_workgroup_2;
   sycl::range<3> max_global_workgroup_3;
+};
 
-  template <int N> inline sycl::range<N> get_max_global_workgroup();
+template <int N>
+inline sycl::range<N> get_max_global_workgroup(WorkGroupLimits &wgl);
 
-  template <> inline sycl::range<1> get_max_global_workgroup() {
-    return this->max_global_workgroup_1;
-  }
-  template <> inline sycl::range<2> get_max_global_workgroup() {
-    return this->max_global_workgroup_2;
-  }
-  template <> inline sycl::range<3> get_max_global_workgroup() {
-    return this->max_global_workgroup_3;
-  }
+template <>
+inline sycl::range<1> get_max_global_workgroup(WorkGroupLimits &wgl) {
+  return wgl.max_global_workgroup_1;
+}
+template <>
+inline sycl::range<2> get_max_global_workgroup(WorkGroupLimits &wgl) {
+  return wgl.max_global_workgroup_2;
+}
+template <>
+inline sycl::range<3> get_max_global_workgroup(WorkGroupLimits &wgl) {
+  return wgl.max_global_workgroup_3;
+}
+
+} // namespace Private
+
+class DeviceLimits {
+protected:
+  Private::WorkGroupLimits wgl;
 
   inline void setup_env() {
-    auto current_limits = this->get_max_global_workgroup<3>();
+    auto current_limits = Private::get_max_global_workgroup<3>(wgl);
     const std::size_t env_max_0 = get_env_size_t(
         "NESO_PARTICLES_DEVICE_LIMIT_GLOBAL_SIZE_S0", current_limits.get(2));
     const std::size_t env_max_1 = get_env_size_t(
         "NESO_PARTICLES_DEVICE_LIMIT_GLOBAL_SIZE_S1", current_limits.get(1));
     const std::size_t env_max_2 = get_env_size_t(
         "NESO_PARTICLES_DEVICE_LIMIT_GLOBAL_SIZE_S2", current_limits.get(0));
 
-    this->max_global_workgroup_1 = sycl::range<1>(env_max_0);
-    this->max_global_workgroup_2 = sycl::range<2>(env_max_1, env_max_0);
-    this->max_global_workgroup_3 =
+    this->wgl.max_global_workgroup_1 = sycl::range<1>(env_max_0);
+    this->wgl.max_global_workgroup_2 = sycl::range<2>(env_max_1, env_max_0);
+    this->wgl.max_global_workgroup_3 =
         sycl::range<3>(env_max_2, env_max_1, env_max_0);
   }
 
   inline void setup_generic() {
     constexpr std::size_t max_size_t = std::numeric_limits<std::size_t>::max();
-    this->max_global_workgroup_1 = sycl::range<1>(max_size_t);
-    this->max_global_workgroup_2 = sycl::range<2>(max_size_t, max_size_t);
-    this->max_global_workgroup_3 =
+    this->wgl.max_global_workgroup_1 = sycl::range<1>(max_size_t);
+    this->wgl.max_global_workgroup_2 = sycl::range<2>(max_size_t, max_size_t);
+    this->wgl.max_global_workgroup_3 =
         sycl::range<3>(max_size_t, max_size_t, max_size_t);
   }
 
   inline void setup_nvidia() {
-    this->max_global_workgroup_1 = sycl::range<1>(2147483647);
-    this->max_global_workgroup_2 = sycl::range<2>(65535, 2147483647);
-    this->max_global_workgroup_3 = sycl::range<3>(65535, 65535, 2147483647);
+    this->wgl.max_global_workgroup_1 = sycl::range<1>(2147483647);
+    this->wgl.max_global_workgroup_2 = sycl::range<2>(65535, 2147483647);
+    this->wgl.max_global_workgroup_3 = sycl::range<3>(65535, 65535, 2147483647);
   }
 
 public:
@@ -68,11 +80,11 @@ class DeviceLimits {
 
   inline void print() {
     nprint("Using global workgroup limits:");
-    auto d1 = this->get_max_global_workgroup<1>();
+    auto d1 = Private::get_max_global_workgroup<1>(wgl);
     nprint("1D:", d1.get(0));
-    auto d2 = this->get_max_global_workgroup<2>();
+    auto d2 = Private::get_max_global_workgroup<2>(wgl);
     nprint("2D:", d2.get(0), d2.get(1));
-    auto d3 = this->get_max_global_workgroup<3>();
+    auto d3 = Private::get_max_global_workgroup<3>(wgl);
     nprint("3D:", d3.get(0), d3.get(1), d3.get(2));
   }
 
@@ -85,7 +97,8 @@ class DeviceLimits {
   template <int N>
   inline sycl::range<N>
   validate_range_global(const sycl::range<N> &range_global) {
-    sycl::range<N> max_global_workgroup = this->get_max_global_workgroup<N>();
+    sycl::range<N> max_global_workgroup =
+        Private::get_max_global_workgroup<N>(wgl);
     for (int dx = 0; dx < N; dx++) {
       if (max_global_workgroup.get(dx)) {
         NESOASSERT(

include/neso_particles/loop/particle_loop.hpp

@@ -427,7 +427,7 @@ class ParticleLoop : public ParticleLoopBase {
     this->particle_dat_init = std::static_pointer_cast<void>(particle_dat);
     this->unpack_args<0>(args...);
     (check_is_sym_outer(args), ...);
-  };
+  }
 
   /**
    *  Create a ParticleLoop that executes a kernel for all particles in the
@@ -442,7 +442,7 @@ class ParticleLoop : public ParticleLoopBase {
   template <typename DAT_TYPE>
   ParticleLoop(ParticleDatSharedPtr<DAT_TYPE> particle_dat, KERNEL kernel,
                ARGS... args)
-      : ParticleLoop("unnamed_kernel", particle_dat, kernel, args...){};
+      : ParticleLoop("unnamed_kernel", particle_dat, kernel, args...) {}
 
   /**
    *  Launch the ParticleLoop and return. Must be called collectively over the

include/neso_particles/particle_group.hpp

@@ -103,15 +103,15 @@ class ParticleGroup {
   template <typename T>
   inline void realloc_dat_start(ParticleDatSharedPtr<T> &dat) {
     dat->realloc(this->h_npart_cell);
-  };
+  }
   template <typename T>
   inline void realloc_dat_wait(ParticleDatSharedPtr<T> &dat) {
     dat->wait_realloc();
-  };
+  }
   template <typename T> inline void realloc_dat(ParticleDatSharedPtr<T> &dat) {
     this->realloc_dat_start(dat);
     this->realloc_dat_wait(dat);
-  };
+  }
 
   inline void realloc_all_dats() {
     for (auto &dat : this->particle_dats_real) {
@@ -207,7 +207,7 @@ class ParticleGroup {
 
   template <typename T> inline void push_particle_spec(ParticleProp<T> prop) {
     this->particle_spec.push(prop);
-  };
+  }
 
   // members for mpi communication
   // global communication context

include/neso_particles/particle_io.hpp

@@ -230,7 +230,7 @@ class H5Part {
                               H5P_DEFAULT, this->plist_id);
     this->is_closed = false;
     this->step = 0;
-  };
+  }
 
   /**
    *  Close the H5Part writer. Must be called before execution completes. Must

include/neso_particles/particle_spec.hpp

@@ -117,7 +117,7 @@ class ParticleSpec {
    *
    *  @param args ParticleSpec is called with a set of ParticleProp arguments.
    */
-  template <typename... T> ParticleSpec(T... args) { this->push(args...); };
+  template <typename... T> ParticleSpec(T... args) { this->push(args...); }
 
   /**
    * Create a particle specification from a vector of REAL properties and a
@@ -218,7 +218,7 @@ class SymStore {
    *
    *  @param args Passed arguments should be Sym<REAL> or Sym<INT>.
    */
-  template <typename... T> SymStore(T... args) { this->push(args...); };
+  template <typename... T> SymStore(T... args) { this->push(args...); }
 
   SymStore() {};
   ~SymStore() {};

include/neso_particles/typedefs.hpp

@@ -12,7 +12,7 @@
 
 namespace NESO::Particles {
 
-static inline int reduce_mul(const int nel, std::vector<int> &values) {
+inline int reduce_mul(const int nel, std::vector<int> &values) {
   int v = 1;
   for (int ex = 0; ex < nel; ex++) {
     v *= values[ex];
@@ -203,13 +203,6 @@ template <typename... T> inline void nprint(T... args) {
 #define NESO_PARTICLES_DEVICE_LABEL "CPU"
 #define NESO_PARTICLES_ITER_CELLS 1
 
-//#define NESO_PARTICLES_KERNEL_START                                            \
-//  const int neso_npart = pl_npart_cell[idx];                                   \
-//  for (int neso_layer = 0; neso_layer < neso_npart; neso_layer++) {
-// #define NESO_PARTICLES_KERNEL_END }
-// #define NESO_PARTICLES_KERNEL_CELL idx
-// #define NESO_PARTICLES_KERNEL_LAYER neso_layer
-
 #define NESO_PARTICLES_KERNEL_START                                            \
   const int neso_cell = (((INT)idx) / pl_stride);                              \
   const int neso_npart = pl_npart_cell[neso_cell];                             \