++

mkn.sh

@@ -8,8 +8,11 @@ cls
 set -e
 
 TEST="-M tests/core/numerics/ion_updater/test_multi_updater.cpp"
-ARGS="${TEST} -P mkn.base=gpu_ -x "
-[ -d /opt/rocm/bin ] && ARGS+="res/mkn/hip" || ARGS+="res/mkn/clang_cuda "
+XFILE="${XFILE:-res/mkn/clang_cuda}"
+[ -d /opt/rocm/bin ] && XFILE="res/mkn/hip"
+
+ARGS="${TEST} -P mkn.base=gpu_"
+[ -n "XFILE" ] && ARGS+=" -x ${XFILE}"
 
 set -x
 
@@ -22,6 +25,3 @@ set -x
   mkn clean build -p test_core ${ARGS} $@
 
 ) #1> >(tee $CWD/.mkn.sh.out ) 2> >(tee $CWD/.mkn.sh.err >&2 )
-
-exit 0 # comment out to do soak test
-

src/core/data/particles/arrays/particle_array_pc.hpp

@@ -432,6 +432,10 @@ class PerCellVector
     {
         on_box_list(local_box().remove(shrink(local_box(box()), 1)), fn);
     };
+    void on_ghost_layer_plus_2_domain(auto&& fn) const
+    {
+        on_box_list(local_box().remove(shrink(local_box(box()), 2)), fn);
+    };
 
 }; // PerCellVector<Particles>
 
@@ -514,7 +518,7 @@ auto& PerCellVector<Particles, impl>::reserve_ppc(std::size_t const& ppc)
 
     if constexpr (type == ParticleType::Ghost)
     {
-        on_ghost_layer_plus_1_domain([&](auto const& bix) {
+        on_ghost_layer_plus_2_domain([&](auto const& bix) {
             particles_(bix).reserve(additional);
             reserve(gaps_(bix), additional);
         });
@@ -683,21 +687,31 @@ void PerCellVector<Particles, impl>::reset_index_wrapper_map()
 {
     resize(p2c_, total_size);
 
+    auto const fill = [](auto p, auto o, auto s, auto b) { std::fill(p + o, p + o + s, *b); };
 
     std::size_t offset = 0;
     for (auto const& bix : local_box())
     {
         auto const& cs = cell_size_(bix);
-        if (cs == 0)
-            continue;
-
         resize(gaps_(bix), cs);
         off_sets_(bix) = offset;
 
-        if constexpr (alloc_mode == AllocatorMode::GPU_UNIFIED)
-            thrust::fill(thrust::device, p2c_.begin() + offset, p2c_.begin() + offset + cs, *bix);
-        else
-            std::fill(p2c_.begin() + offset, p2c_.begin() + offset + cs, *bix);
+        if (cs)
+        {
+            if constexpr (alloc_mode == AllocatorMode::GPU_UNIFIED)
+            {
+                PHARE_WITH_THRUST( //
+                    thrust::fill(thrust::device, p2c_.begin() + offset, p2c_.begin() + offset + cs,
+                                 *bix));
+                PHARE_WITH_THRUST_ELSE(
+                    PHARE_LOG_LINE_SS("Thrust not found for PerCellVector<Particles, "
+                                      "impl>::reset_index_wrapper_map"); //
+                    fill(p2c_.begin(), offset, cs, bix);                 //
+                )
+            }
+            else
+                fill(p2c_.begin(), offset, cs, bix);
+        }
 
         offset += cs;
         cap_(bix) = particles_(bix).capacity();
@@ -867,11 +881,12 @@ struct PerCellParticles : public Super_
             using Op = Operators<typename Super::SIZE_T, true>;
 
             // printf("L:%d i %llu ic %u,%u change \n", __LINE__, idx, cell[0], cell[1]);
+            Super::gaps_(cell)[Op{Super::gap_idx_(cell)}.increment_return_old()] = idx;
+
             if (isIn(newcell, Super::ghost_box()))
             {
                 auto const nc = Super::local_cell(newcell);
                 Op{Super::add_into_(nc)}.increment_return_old();
-                Super::gaps_(cell)[Op{Super::gap_idx_(cell)}.increment_return_old()] = idx;
             }
         }
         else
@@ -1005,25 +1020,30 @@ struct PerCellParticles<OuterSuper>::iterator_impl
 template<auto layout_mode, typename Particles>
 struct index_wrapper_storage;
 
+#if PHARE_HAVE_THRUST
 template<typename Particles>
 struct index_wrapper_storage<LayoutMode::SoA, Particles>
 {
     bool static constexpr is_const = std::is_const_v<std::remove_reference_t<Particles>>;
     using per_cell_particles       = typename std::decay_t<Particles>::per_cell_particles;
-    using Particle_t               = detail::SoAZipParticle<per_cell_particles>;
+    using Particle_t = typename SoAZipParticle_t<per_cell_particles, is_const>::value_type;
 
-    index_wrapper_storage(per_cell_particles* p, std::size_t const i) _PHARE_ALL_FN_
-        : particles{p},
+    template<typename PerCellParticles_t>
+    index_wrapper_storage(PerCellParticles_t p, std::size_t const i) _PHARE_ALL_FN_
+        : /*particles{p},*/
           particle{*p, i}
     {
     }
 
-    auto& deref() _PHARE_ALL_FN_ { return particle; }
-    // auto& reset(std::size_t const i) { return particle = detail::SoAZipParticle(*particles, i); }
+    auto& operator*() _PHARE_ALL_FN_ { return particle; }
+    auto& operator*() const _PHARE_ALL_FN_ { return particle; }
 
     per_cell_particles* particles;
     Particle_t particle;
 };
+#else
+
+#endif // PHARE_HAVE_THRUST
 
 template<typename Particles>
 struct index_wrapper_storage<LayoutMode::AoS, Particles>
@@ -1033,16 +1053,17 @@ struct index_wrapper_storage<LayoutMode::AoS, Particles>
     using Particle_t               = typename Particles::Particle_t;
     using Particle_p = std::conditional_t<is_const, Particle_t const* const, Particle_t*>;
 
-    index_wrapper_storage(per_cell_particles* p, std::size_t const i) _PHARE_ALL_FN_
-        : particles{p},
+    template<typename PerCellParticles_t>
+    index_wrapper_storage(PerCellParticles_t p, std::size_t const i) _PHARE_ALL_FN_
+        : /*particles{p},*/
           particle{&p->data()[i]}
     {
     }
 
-    auto& deref() _PHARE_ALL_FN_ { return *particle; }
-    // auto& reset(std::size_t const i) { return particle = p.data()[i]; }
+    auto& operator*() _PHARE_ALL_FN_ { return *particle; }
+    auto& operator*() const _PHARE_ALL_FN_ { return *particle; }
 
-    per_cell_particles* particles;
+    // per_cell_particles* particles;
     Particle_p particle;
 };
 
@@ -1065,7 +1086,6 @@ struct PerCellParticles<ParticlesSuper>::index_wrapper : public index_wrapper_su
 {
     using outer_t = std::decay_t<T>;
     using Super   = typename index_wrapper_super<T>::value_type;
-    using Super::deref;
 
     auto static constexpr dimension = ParticlesSuper::dimension;
     bool static constexpr is_const  = std::is_const_v<std::remove_reference_t<T>>;
@@ -1078,11 +1098,11 @@ struct PerCellParticles<ParticlesSuper>::index_wrapper : public index_wrapper_su
           pc_particles_ptr{pc_particles},
           idx{idx_}
     {
-        PHARE_ASSERT(deref().iCell()[0] > -10 and deref().iCell()[0] < 1000); // bad memory
+        PHARE_ASSERT((**this).iCell()[0] > -10 and (**this).iCell()[0] < 1000); // bad memory
         if constexpr (dimension > 1)
-            PHARE_ASSERT(deref().iCell()[1] > -10 and deref().iCell()[1] < 1000); // bad memory
+            PHARE_ASSERT((**this).iCell()[1] > -10 and (**this).iCell()[1] < 1000); // bad memory
         if constexpr (dimension > 2)
-            PHARE_ASSERT(deref().iCell()[2] > -10 and deref().iCell()[2] < 1000); // bad memory
+            PHARE_ASSERT((**this).iCell()[2] > -10 and (**this).iCell()[2] < 1000); // bad memory
     }
 
     auto& c() const _PHARE_ALL_FN_ { return cell(pc_particles_ptr, idx); }
@@ -1096,7 +1116,20 @@ struct PerCellParticles<ParticlesSuper>::index_wrapper : public index_wrapper_su
 
     auto icell_changer(std::array<int, dimension> const& newcell) _PHARE_ALL_FN_
     {
-        pc_particles_ptr->icell_changer(deref(), c(), i(), newcell);
+        pc_particles_ptr->icell_changer(**this, c(), i(), newcell);
+    }
+
+
+    Super& super() _PHARE_ALL_FN_ { return *this; }
+    Super const& super() const _PHARE_ALL_FN_ { return *this; }
+
+    auto& operator*() _PHARE_ALL_FN_ { return *super(); }
+    auto& operator*() const _PHARE_ALL_FN_ { return *super(); }
+
+    Particle<dimension> copy() const _PHARE_ALL_FN_
+    {
+        return {(**this).weight(), (**this).charge(), (**this).iCell(), (**this).delta(),
+                (**this).v()};
     }
 
 

src/core/data/particles/arrays/particle_array_soa.hpp

@@ -518,11 +518,11 @@ class SoAParticles : public Super_
 
     auto operator[](std::size_t const& s) const _PHARE_ALL_FN_
     {
-#if __has_include(<thrust/iterator/zip_iterator.h>)
-        return detail::SoAZipConstParticle(*this, s);
-#else
+        // #if __has_include(<thrust/iterator/zip_iterator.h>)
+        //         return detail::SoAZipConstParticle(*this, s);
+        // #else
         return copy(s);
-#endif // __has_include(<thrust/iterator/zip_iterator.h>)
+        // #endif // __has_include(<thrust/iterator/zip_iterator.h>)
     }
 };
 

src/core/data/particles/arrays/particle_array_soa_thrust.hpp

@@ -156,40 +156,48 @@ struct SoAZipConstParticle
         std::declval<SoAParticles_t&>()))>;
 
     SoAZipConstParticle(SoAParticles_t& ps, std::size_t const& i) _PHARE_ALL_FN_
-        : it{SoAIteratorAdaptor::make(ps, i)},
-          ref{weight(), charge(), iCell(), delta(), v()}
+        : it{SoAIteratorAdaptor::make(ps, i)} /*,
+           ref{weight(), charge(), iCell(), delta(), v()}*/
     {
     }
 
-    auto& charge() _PHARE_ALL_FN_ { return SoAIteratorAdaptor::charge(*it); }
+
     auto& charge() const _PHARE_ALL_FN_ { return SoAIteratorAdaptor::charge(*it); }
-    auto& weight() _PHARE_ALL_FN_ { return SoAIteratorAdaptor::weight(*it); }
     auto& weight() const _PHARE_ALL_FN_ { return SoAIteratorAdaptor::weight(*it); }
-
-    auto& iCell() _PHARE_ALL_FN_ { return SoAIteratorAdaptor::iCell(*it); }
     auto& iCell() const _PHARE_ALL_FN_ { return SoAIteratorAdaptor::iCell(*it); }
-    auto& delta() _PHARE_ALL_FN_ { return SoAIteratorAdaptor::delta(*it); }
     auto& delta() const _PHARE_ALL_FN_ { return SoAIteratorAdaptor::delta(*it); }
-
-
-    auto& v() _PHARE_ALL_FN_ { return SoAIteratorAdaptor::v(*it); }
     auto& v() const _PHARE_ALL_FN_ { return SoAIteratorAdaptor::v(*it); }
 
-    auto& operator*() _PHARE_ALL_FN_ { return ref; }
-    auto& operator*() const _PHARE_ALL_FN_ { return ref; }
+    // auto& operator*() _PHARE_ALL_FN_ { return ref; }
+    // auto& operator*() const _PHARE_ALL_FN_ { return ref; }
 
     Iterator it;
-    SoAParticle_crt<dim> ref;
+    // SoAParticle_crt<dim> ref;
 };
 
 
-
 } // namespace PHARE::core::detail
 
 
 namespace PHARE::core
 {
 
+template<typename SoAParticles_t, bool _is_const = false>
+struct SoAZipParticle_t
+{
+    bool static constexpr is_const
+        = _is_const || std::is_const_v<std::remove_reference_t<SoAParticles_t>>;
+
+    using value_type = std::conditional_t<is_const, detail::SoAZipConstParticle<SoAParticles_t>,
+                                          detail::SoAZipParticle<SoAParticles_t>>;
+};
+
+template<typename Particles>
+auto particle_zip_iterator(Particles& ps, std::size_t const i)
+{
+    return typename SoAZipParticle_t<Particles>::value_type{ps, i};
+}
+
 template<typename T, std::size_t dim>
 auto partitionner(detail::SoAIteratorAdaptor& begin, detail::SoAIteratorAdaptor& end,
                   Box<T, dim> const& box)

src/core/def/thrust.hpp

@@ -3,23 +3,28 @@
 
 #if __has_include(<thrust/iterator/zip_iterator.h>)
 
-#include <thrust/iterator/zip_iterator.h>
 #define PHARE_HAVE_THRUST 1
 #define PHARE_WITH_THRUST(...) __VA_ARGS__
+#define PHARE_WITH_THRUST_ELSE(...)
 #define PHARE_WITH_THRUST_ELSE_THROW(...) __VA_ARGS__
 
-#else // !__has_include(...)
+#else // !__has_include(<thrust/iterator/zip_iterator.h>)
 
 
 #define PHARE_HAVE_THRUST 0
 #define PHARE_WITH_THRUST(...)
+#define PHARE_WITH_THRUST_ELSE(...) __VA_ARGS__
 #define PHARE_WITH_THRUST_ELSE_THROW(...) throw std::runtime_error("Thrust not found!");
 
-#endif // __has_include(...)
+#endif // __has_include(<thrust/iterator/zip_iterator.h>)
+
 
 #if PHARE_HAVE_THRUST
+
 #include <thrust/sort.h>
 #include <thrust/execution_policy.h>
+#include <thrust/iterator/zip_iterator.h>
+
 #endif // PHARE_HAVE_THRUST
 
 #endif /* PHARE_CORE_DEF_THRUST_HPP */

src/core/numerics/interpolator/interpolating.hpp

@@ -49,8 +49,8 @@ class Interpolating
                 static_assert(atomic_ops, "GPU must be atomic");
                 PHARE_WITH_MKN_GPU(
                     mkn::gpu::GDLauncher{particles.size()}([=] _PHARE_ALL_FN_() mutable {
-                        Interpolator_t{}.particleToMesh(particles[mkn::gpu::idx()], density, flux,
-                                                        layout, coef);
+                        auto it = particles[mkn::gpu::idx()];
+                        Interpolator_t{}.particleToMesh(*it, density, flux, layout, coef);
                     }); //
                 )
             }

src/core/numerics/pusher/boris_simpler.hpp

@@ -276,13 +276,14 @@ class SimpleBorisPusher
             auto per_particle = [=] _PHARE_ALL_FN_<bool accelerate = false>() mutable
             {
                 Interpolator interp;
-                auto particle = view[mkn::gpu::idx()];
+                auto it        = view[mkn::gpu::idx()];
+                auto& particle = *it;
                 if constexpr (accelerate)
                     boris_accelerate(particle, interp.m2p(particle, em, layout), dto2m_);
                 auto const& newCell = advancePosition_<alloc_mode>(particle, halfDtOverDl);
                 if (!array_equals(newCell, particle.iCell()))
                 {
-                    particle.icell_changer(newCell);
+                    it.icell_changer(newCell);
                     particle.iCell() = newCell;
                 }
             };