Improve single-event mode in accel

example/accel/fastsim-offload.cc

@@ -209,8 +209,6 @@ int main()
     // NOTE: these numbers are appropriate for CPU execution
     setup_options.max_num_tracks = 1024;
     setup_options.initializer_capacity = 1024 * 128;
-    // This parameter will eventually be removed
-    setup_options.max_num_events = 1024;
     // Celeritas does not support EmStandard MSC physics above 100 MeV
     setup_options.ignore_processes = {"CoulombScat"};
     if (G4VERSION_NUMBER >= 1110)

src/accel/SetupOptions.hh

@@ -119,7 +119,7 @@ struct SetupOptions
     //! \name Celeritas stepper options
     //! Number of track "slots" to be transported simultaneously
     size_type max_num_tracks{};
-    //! Maximum number of events in use
+    //! Maximum number of events in use (DEPRECATED: remove in v0.6)
     size_type max_num_events{};
     //! Limit on number of step iterations before aborting
     size_type max_steps = no_max_steps();

src/accel/SetupOptionsMessenger.cc

@@ -180,7 +180,7 @@ SetupOptionsMessenger::SetupOptionsMessenger(SetupOptions* options)
             "Number of track \"slots\" to be transported simultaneously");
     add_cmd(&options->max_num_events,
             "maxNumEvents",
-            "Maximum number of events in use");
+            "Maximum number of events in use (DEPRECATED)");
     add_cmd(&options->max_steps,
             "maxNumSteps",
             "Limit on number of step iterations before aborting");

src/accel/SharedParams.cc

@@ -530,11 +530,17 @@ void SharedParams::initialize_core(SetupOptions const& options)
     params.sim = SimParams::from_import(
         *imported, params.particle, options.max_field_substeps);
 
+    if (options.max_num_events > 0)
+    {
+        CELER_LOG(warning) << "Deprecated option 'max_events': will be "
+                              "removed in v0.6";
+    }
+
     // Construct track initialization params
     params.init = [&options] {
         TrackInitParams::Input input;
         input.capacity = options.initializer_capacity;
-        input.max_events = options.max_num_events;
+        input.max_events = 1;  // TODO: use special "max events" case
         input.track_order = options.track_order;
         return std::make_shared<TrackInitParams>(std::move(input));
     }();

src/celeritas/CMakeLists.txt

@@ -125,7 +125,7 @@ endif()
 
 if(CELERITAS_USE_CUDA OR CELERITAS_USE_HIP)
   list(APPEND SOURCES
-    track/detail/Filler.cu
+    track/Filler.cu
     global/alongstep/detail/AlongStepKernels.cu
   )
 endif()

src/celeritas/global/Stepper.cc

@@ -165,7 +165,7 @@ auto Stepper<M>::operator()(SpanConstPrimary primaries) -> result_type
 
 //---------------------------------------------------------------------------//
 /*!
- * Reseed the RNGs at the start of an event for "strong" reproducibility.
+ * Reseed RNGs and counters the start of an event for "strong" reproducibility.
  *
  * This reinitializes the RNG states using a single seed and unique subsequence
  * for each thread. It ensures that given an event number, that event can be
@@ -174,7 +174,11 @@ auto Stepper<M>::operator()(SpanConstPrimary primaries) -> result_type
 template<MemSpace M>
 void Stepper<M>::reseed(UniqueEventId event_id)
 {
-    reseed_rng(get_ref<M>(*params_->rng()), state_->ref().rng, event_id);
+    reseed_rng(get_ref<M>(*params_->rng()),
+               state_->ref().rng,
+               state_->stream_id(),
+               event_id);
+    params_->init()->reset_track_ids(state_->stream_id(), &state_->ref().init);
 }
 
 //---------------------------------------------------------------------------//

src/celeritas/random/RngReseed.cc

@@ -24,6 +24,7 @@ namespace celeritas
  */
 void reseed_rng(HostCRef<RngParamsData> const& params,
                 HostRef<RngStateData> const& state,
+                StreamId,
                 UniqueEventId event_id)
 {
     CELER_EXPECT(event_id);

src/celeritas/random/RngReseed.cu

@@ -57,12 +57,19 @@ __global__ void reseed_rng_kernel(DeviceCRef<RngParamsData> const params,
  */
 void reseed_rng(DeviceCRef<RngParamsData> const& params,
                 DeviceRef<RngStateData> const& state,
+                StreamId stream,
                 UniqueEventId event_id)
 {
     CELER_EXPECT(state);
     CELER_EXPECT(params);
-    CELER_LAUNCH_KERNEL(
-        reseed_rng, state.size(), 0, params, state, event_id.get());
+    CELER_EXPECT(stream);
+
+    CELER_LAUNCH_KERNEL(reseed_rng,
+                        state.size(),
+                        celeritas::device().stream(state.stream_id).get(),
+                        params,
+                        state,
+                        event_id.get());
 }
 
 //---------------------------------------------------------------------------//

src/celeritas/random/RngReseed.hh

@@ -21,10 +21,12 @@ namespace celeritas
 // Reinitialize the RNG states on host/device at the start of an event
 void reseed_rng(DeviceCRef<RngParamsData> const&,
                 DeviceRef<RngStateData> const&,
+                StreamId,
                 UniqueEventId);
 
 void reseed_rng(HostCRef<RngParamsData> const&,
                 HostRef<RngStateData> const&,
+                StreamId,
                 UniqueEventId);
 
 #if !CELER_USE_DEVICE
@@ -34,6 +36,7 @@ void reseed_rng(HostCRef<RngParamsData> const&,
  */
 inline void reseed_rng(DeviceCRef<RngParamsData> const&,
                        DeviceRef<RngStateData> const&,
+                       StreamId,
                        UniqueEventId)
 {
     CELER_ASSERT_UNREACHABLE();

src/celeritas/track/TrackInitParams.cc

@@ -10,6 +10,7 @@
 #include <utility>
 
 #include "corecel/Assert.hh"
+#include "corecel/data/Filler.hh"
 
 #include "TrackInitData.hh"  // IWYU pragma: associated
 

src/celeritas/track/TrackInitParams.hh

@@ -9,6 +9,7 @@
 
 #include "corecel/Types.hh"
 #include "corecel/data/CollectionMirror.hh"
+#include "corecel/data/Filler.hh"
 #include "corecel/data/ParamsDataInterface.hh"
 
 #include "TrackInitData.hh"
@@ -26,7 +27,7 @@ class TrackInitParams final : public ParamsDataInterface<TrackInitParamsData>
     struct Input
     {
         size_type capacity{};  //!< Max number of initializers
-        size_type max_events{};  //!< Max number of events that can be run
+        size_type max_events{};  //!< Max simultaneous events
         TrackOrder track_order{TrackOrder::none};  //!< How to sort tracks
     };
 
@@ -49,10 +50,30 @@ class TrackInitParams final : public ParamsDataInterface<TrackInitParamsData>
     //! Access data on device
     DeviceRef const& device_ref() const final { return data_.device_ref(); }
 
+    // Reset track counters in single-event mode
+    template<MemSpace M>
+    inline void
+    reset_track_ids(StreamId stream,
+                    TrackInitStateData<Ownership::reference, M>* state) const;
+
   private:
     // Host/device storage and reference
     CollectionMirror<TrackInitParamsData> data_;
 };
 
+//---------------------------------------------------------------------------//
+/*!
+ * Reset track counters in single-event mode.
+ */
+template<MemSpace M>
+void TrackInitParams::reset_track_ids(
+    StreamId stream, TrackInitStateData<Ownership::reference, M>* state) const
+{
+    CELER_EXPECT(state);
+
+    Filler<size_type, M> fill_zero{0, stream};
+    fill_zero(state->track_counters[AllItems<size_type, M>{}]);
+}
+
 //---------------------------------------------------------------------------//
 }  // namespace celeritas

src/corecel/data/CollectionAlgorithms.hh

@@ -12,22 +12,24 @@
 
 #include "Collection.hh"
 #include "Copier.hh"
-
-#include "detail/Filler.hh"
+#include "Filler.hh"
 
 namespace celeritas
 {
 //---------------------------------------------------------------------------//
 /*!
  * Fill the collection with the given value.
+ *
+ * This should only be used during initialization when stream synchronization
+ * is OK. Use the \c Filler directly during runtime since it supports streams.
  */
 template<class T, Ownership W, MemSpace M, class I>
 void fill(T&& value, Collection<T, W, M, I>* col)
 {
     static_assert(W != Ownership::const_reference,
                   "const references cannot be filled");
     CELER_EXPECT(col);
-    detail::Filler<T, M> fill_impl{value};
+    Filler<T, M> fill_impl{value};
     fill_impl((*col)[AllItems<T, M>{}]);
 }
 

src/corecel/data/Copier.hh

@@ -38,10 +38,10 @@ class Copier
 
   public:
     //! Construct with the destination and the class's memspace
-    explicit Copier(Span<T> dst) : dst_{dst} {};
+    explicit Copier(Span<T> dst) : dst_{dst} {}
 
     //! Also construct with a stream ID to use for async copy
-    Copier(Span<T> dst, StreamId stream) : dst_{dst}, stream_{stream} {};
+    Copier(Span<T> dst, StreamId stream) : dst_{dst}, stream_{stream} {}
 
     inline void operator()(MemSpace srcmem, Span<T const> src) const;
 
@@ -53,7 +53,7 @@ class Copier
 
 //---------------------------------------------------------------------------//
 /*!
- * Copy a value from device to host.
+ * Copy a single value from device to host.
  *
  * The source of the data to copy is the function argument.
  */
@@ -68,7 +68,7 @@ class ItemCopier
     ItemCopier() = default;
 
     //! Also construct with a stream ID to use for async copy
-    explicit ItemCopier(StreamId stream) : stream_{stream} {};
+    explicit ItemCopier(StreamId stream) : stream_{stream} {}
 
     inline T operator()(T const* src) const;
 

src/corecel/data/detail/Filler.cu → src/corecel/data/Filler.cu

@@ -3,19 +3,16 @@
 // See the top-level COPYRIGHT file for details.
 // SPDX-License-Identifier: (Apache-2.0 OR MIT)
 //---------------------------------------------------------------------------//
-//! \file corecel/data/detail/Filler.cu
+//! \file corecel/data/Filler.cu
 //---------------------------------------------------------------------------//
 #include "Filler.device.t.hh"
 
 namespace celeritas
 {
-namespace detail
-{
 //---------------------------------------------------------------------------//
 template struct Filler<real_type, MemSpace::device>;
 template struct Filler<size_type, MemSpace::device>;
 template struct Filler<int, MemSpace::device>;
 template struct Filler<TrackSlotId, MemSpace::device>;
 //---------------------------------------------------------------------------//
-}  // namespace detail
 }  // namespace celeritas

src/corecel/data/detail/Filler.device.t.hh → src/corecel/data/Filler.device.t.hh

@@ -3,7 +3,7 @@
 // See the top-level COPYRIGHT file for details.
 // SPDX-License-Identifier: (Apache-2.0 OR MIT)
 //---------------------------------------------------------------------------//
-//! \file corecel/data/detail/Filler.device.t.hh
+//! \file corecel/data/Filler.device.t.hh
 //---------------------------------------------------------------------------//
 #pragma once
 
@@ -20,11 +20,9 @@
 
 namespace celeritas
 {
-namespace detail
-{
 //---------------------------------------------------------------------------//
-template<class T>
-void Filler<T, MemSpace::device>::operator()(Span<T> data) const
+template<class T, MemSpace M>
+void Filler<T, M>::fill_device_impl(Span<T>) const
 {
     if (stream_)
     {
@@ -44,5 +42,4 @@ void Filler<T, MemSpace::device>::operator()(Span<T> data) const
 }
 
 //---------------------------------------------------------------------------//
-}  // namespace detail
 }  // namespace celeritas

src/corecel/data/detail/Filler.hh → src/corecel/data/Filler.hh

@@ -3,7 +3,7 @@
 // See the top-level COPYRIGHT file for details.
 // SPDX-License-Identifier: (Apache-2.0 OR MIT)
 //---------------------------------------------------------------------------//
-//! \file corecel/data/detail/Filler.hh
+//! \file corecel/data/Filler.hh
 //---------------------------------------------------------------------------//
 #pragma once
 
@@ -17,43 +17,54 @@
 
 namespace celeritas
 {
-namespace detail
-{
 //---------------------------------------------------------------------------//
-
-//! Assign on host or mapped / managed memory
+/*!
+ * Fill contiguous data with copies of a scalar value.
+ */
 template<class T, MemSpace M>
-struct Filler
-{
-    T const& value;
-
-    void operator()(Span<T> data) const
-    {
-        std::fill(data.begin(), data.end(), value);
-    }
-};
-
-//! Assign on device
-template<class T>
-struct Filler<T, MemSpace::device>
+class Filler
 {
   public:
-    explicit Filler(T const& value) : value_{value} {};
-    Filler(T const& value, StreamId stream)
-        : value_{value}, stream_{stream} {};
+    //! Construct with target value and default stream
+    explicit Filler(T value) : value_{value} {}
+
+    //! Construct with target value and specific stream
+    Filler(T value, StreamId stream) : value_{value}, stream_{stream} {}
 
-    void operator()(Span<T>) const;
+    // Fill the span with the stored value
+    inline void operator()(Span<T> data) const;
 
   private:
-    T const& value_;
+    T value_;
     StreamId stream_;
+
+    void fill_device_impl(Span<T> data) const;
 };
 
+//---------------------------------------------------------------------------//
+/*!
+ * Fill the span with the stored value.
+ */
+template<class T, MemSpace M>
+void Filler<T, M>::operator()(Span<T> data) const
+{
+    if constexpr (M == MemSpace::device)
+    {
+        this->fill_device_impl(data);
+    }
+    else
+    {
+        std::fill(data.begin(), data.end(), value_);
+    }
+}
+
+//---------------------------------------------------------------------------//
 #if !CELER_USE_DEVICE
-template<class T>
-void Filler<T, MemSpace::device>::operator()(Span<T>) const
+template<class T, MemSpace M>
+CELER_FORCEINLINE void Filler<T, M>::fill_device_impl(Span<T>) const
 {
-    CELER_NOT_CONFIGURED("CUDA or HIP");
+    CELER_DISCARD(stream_);
+    CELER_ASSERT_UNREACHABLE();
 }
 #else
 extern template struct Filler<real_type, MemSpace::device>;
@@ -63,5 +74,4 @@ extern template struct Filler<TrackSlotId, MemSpace::device>;
 #endif
 
 //---------------------------------------------------------------------------//
-}  // namespace detail
 }  // namespace celeritas

test/celeritas/geo/HeuristicGeoTestBase.cc

@@ -11,7 +11,6 @@
 #include <iostream>
 
 #include "corecel/cont/Range.hh"
-#include "corecel/data/CollectionAlgorithms.hh"
 #include "corecel/data/CollectionStateStore.hh"
 #include "corecel/data/Ref.hh"
 #include "corecel/io/Join.hh"