Define a "unique event ID" different from event counter (celeritas-pr…

…oject#1447)

* Use a longer integer for event reseeding

Experiment frameworks expect to have more than 4 billion globally unique
event identifiers, and these aren't run sequentially in a single
simulation. Define a "unique event ID" as UID-like integer only for
reseeding and not for tracking.

* Force event ID to be zero

src/accel/LocalTransporter.cc

@@ -122,7 +122,7 @@ void LocalTransporter::InitializeEvent(int id)
     CELER_EXPECT(*this);
     CELER_EXPECT(id >= 0);
 
-    event_id_ = EventId(id);
+    event_id_ = UniqueEventId(id);
     track_counter_ = 0;
 
     if (!(G4Threading::IsMultithreadedApplication()
@@ -142,7 +142,6 @@ void LocalTransporter::InitializeEvent(int id)
 void LocalTransporter::Push(G4Track const& g4track)
 {
     CELER_EXPECT(*this);
-    CELER_EXPECT(event_id_);
 
     Primary track;
 
@@ -173,7 +172,7 @@ void LocalTransporter::Push(G4Track const& g4track)
     // (and share with sensitive detectors!) a map of track IDs for calling
     // back to Geant4.
     track.track_id = TrackId{track_counter_++};
-    track.event_id = event_id_;
+    track.event_id = EventId{0};
 
     buffer_.push_back(track);
     if (buffer_.size() >= auto_flush_)

src/accel/LocalTransporter.hh

@@ -101,7 +101,7 @@ class LocalTransporter
     std::shared_ptr<StepperInterface> step_;
     std::vector<Primary> buffer_;
 
-    EventId event_id_;
+    UniqueEventId event_id_;
     TrackId::size_type track_counter_{};
 
     size_type auto_flush_{};

src/celeritas/Types.hh

@@ -31,9 +31,12 @@ using ActionId = OpaqueId<class ActionInterface>;
 //! Opaque index to ElementRecord in the global vector of elements
 using ElementId = OpaqueId<struct ElementRecord>;
 
-//! Counter for the initiating event for a track
+//! Zero-indexed counter for the initiating event for a track
 using EventId = OpaqueId<struct Event_>;
 
+//! Unique identifier for an event used by external applications
+using UniqueEventId = OpaqueId<struct Event_, std::uint64_t>;
+
 //! Opaque index to IsotopeRecord in a vector
 using IsotopeId = OpaqueId<struct IsotopeRecord>;
 

src/celeritas/global/Stepper.cc

@@ -114,9 +114,9 @@ auto Stepper<M>::operator()(SpanConstPrimary primaries) -> result_type
  * reproduced.
  */
 template<MemSpace M>
-void Stepper<M>::reseed(EventId event_id)
+void Stepper<M>::reseed(UniqueEventId event_id)
 {
-    reseed_rng(get_ref<M>(*params_->rng()), state_.ref().rng, event_id.get());
+    reseed_rng(get_ref<M>(*params_->rng()), state_.ref().rng, event_id);
 }
 
 //---------------------------------------------------------------------------//

src/celeritas/global/Stepper.hh

@@ -91,7 +91,7 @@ class StepperInterface
     virtual StepperResult operator()(SpanConstPrimary primaries) = 0;
 
     // Reseed the RNGs at the start of an event for reproducibility
-    virtual void reseed(EventId event_id) = 0;
+    virtual void reseed(UniqueEventId event_id) = 0;
 
     //! Get action sequence for timing diagnostics
     virtual ActionSequence const& actions() const = 0;
@@ -143,7 +143,7 @@ class Stepper final : public StepperInterface
     StepperResult operator()(SpanConstPrimary primaries) final;
 
     // Reseed the RNGs at the start of an event for reproducibility
-    void reseed(EventId event_id) final;
+    void reseed(UniqueEventId event_id) final;
 
     //! Get action sequence for timing diagnostics
     ActionSequence const& actions() const final { return *actions_; }

src/celeritas/random/RngReseed.cc

@@ -24,14 +24,18 @@ namespace celeritas
  */
 void reseed_rng(HostCRef<RngParamsData> const& params,
                 HostRef<RngStateData> const& state,
-                size_type event_id)
+                UniqueEventId event_id)
 {
-    for (auto tid : range(TrackSlotId{state.size()}))
+    CELER_EXPECT(event_id);
+    static_assert(sizeof(ull_int) == sizeof(UniqueEventId::size_type));
+
+    auto size = static_cast<ull_int>(state.size());
+    for (auto i : range(size))
     {
         RngEngine::Initializer_t init;
         init.seed = params.seed;
-        init.subsequence = event_id * state.size() + tid.get();
-        RngEngine engine(params, state, tid);
+        init.subsequence = event_id.unchecked_get() * size + i;
+        RngEngine engine(params, state, TrackSlotId{i});
         engine = init;
     }
 }

src/celeritas/random/RngReseed.cu

@@ -26,7 +26,7 @@ namespace
  */
 __global__ void reseed_rng_kernel(DeviceCRef<RngParamsData> const params,
                                   DeviceRef<RngStateData> const state,
-                                  size_type event_id)
+                                  UniqueEventId::size_type event_id)
 {
     auto tid = TrackSlotId{
         celeritas::KernelParamCalculator::thread_id().unchecked_get()};
@@ -56,11 +56,12 @@ __global__ void reseed_rng_kernel(DeviceCRef<RngParamsData> const params,
  */
 void reseed_rng(DeviceCRef<RngParamsData> const& params,
                 DeviceRef<RngStateData> const& state,
-                size_type event_id)
+                UniqueEventId event_id)
 {
     CELER_EXPECT(state);
     CELER_EXPECT(params);
-    CELER_LAUNCH_KERNEL(reseed_rng, state.size(), 0, params, state, event_id);
+    CELER_LAUNCH_KERNEL(
+        reseed_rng, state.size(), 0, params, state, event_id.get());
 }
 
 //---------------------------------------------------------------------------//

src/celeritas/random/RngReseed.hh

@@ -11,6 +11,7 @@
 #include "corecel/Macros.hh"
 #include "corecel/Types.hh"
 #include "corecel/data/Collection.hh"
+#include "celeritas/Types.hh"
 
 #include "RngData.hh"
 
@@ -20,11 +21,11 @@ namespace celeritas
 // Reinitialize the RNG states on host/device at the start of an event
 void reseed_rng(DeviceCRef<RngParamsData> const&,
                 DeviceRef<RngStateData> const&,
-                size_type);
+                UniqueEventId);
 
 void reseed_rng(HostCRef<RngParamsData> const&,
                 HostRef<RngStateData> const&,
-                size_type);
+                UniqueEventId);
 
 #if !CELER_USE_DEVICE
 //---------------------------------------------------------------------------//
@@ -33,7 +34,7 @@ void reseed_rng(HostCRef<RngParamsData> const&,
  */
 inline void reseed_rng(DeviceCRef<RngParamsData> const&,
                        DeviceRef<RngStateData> const&,
-                       size_type)
+                       UniqueEventId)
 {
     CELER_ASSERT_UNREACHABLE();
 }

test/celeritas/random/RngReseed.test.cc

@@ -37,7 +37,7 @@ TEST_F(RngReseedTest, reseed)
     RngHostStore states(params->host_ref(), StreamId{0}, size);
 
     size_type id = 8;
-    reseed_rng(params->host_ref(), states.ref(), id);
+    reseed_rng(params->host_ref(), states.ref(), UniqueEventId{id});
 
     RngEngine::Initializer_t init;
     init.seed = params->host_ref().seed;
@@ -55,7 +55,7 @@ TEST_F(RngReseedTest, reseed)
         ASSERT_EQ(skip_rng(), rng());
     }
 
-    reseed_rng(params->host_ref(), states.ref(), id);
+    reseed_rng(params->host_ref(), states.ref(), UniqueEventId{id});
     std::vector<unsigned int> values;
     for (auto i : range(states.size()).step(128u))
     {