diff --git a/src/accel/LocalTransporter.cc b/src/accel/LocalTransporter.cc index 8bf8d8e0f9..c9f7364667 100644 --- a/src/accel/LocalTransporter.cc +++ b/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_) diff --git a/src/accel/LocalTransporter.hh b/src/accel/LocalTransporter.hh index 89b5d68b10..ed5526586e 100644 --- a/src/accel/LocalTransporter.hh +++ b/src/accel/LocalTransporter.hh @@ -101,7 +101,7 @@ class LocalTransporter std::shared_ptr step_; std::vector buffer_; - EventId event_id_; + UniqueEventId event_id_; TrackId::size_type track_counter_{}; size_type auto_flush_{}; diff --git a/src/celeritas/Types.hh b/src/celeritas/Types.hh index 8b524a2a79..9dd6825fb8 100644 --- a/src/celeritas/Types.hh +++ b/src/celeritas/Types.hh @@ -31,9 +31,12 @@ using ActionId = OpaqueId; //! Opaque index to ElementRecord in the global vector of elements using ElementId = OpaqueId; -//! Counter for the initiating event for a track +//! Zero-indexed counter for the initiating event for a track using EventId = OpaqueId; +//! Unique identifier for an event used by external applications +using UniqueEventId = OpaqueId; + //! Opaque index to IsotopeRecord in a vector using IsotopeId = OpaqueId; diff --git a/src/celeritas/global/Stepper.cc b/src/celeritas/global/Stepper.cc index 865b3906ef..1b8b958425 100644 --- a/src/celeritas/global/Stepper.cc +++ b/src/celeritas/global/Stepper.cc @@ -114,9 +114,9 @@ auto Stepper::operator()(SpanConstPrimary primaries) -> result_type * reproduced. */ template -void Stepper::reseed(EventId event_id) +void Stepper::reseed(UniqueEventId event_id) { - reseed_rng(get_ref(*params_->rng()), state_.ref().rng, event_id.get()); + reseed_rng(get_ref(*params_->rng()), state_.ref().rng, event_id); } //---------------------------------------------------------------------------// diff --git a/src/celeritas/global/Stepper.hh b/src/celeritas/global/Stepper.hh index ae3c8a5c6d..a43759fc80 100644 --- a/src/celeritas/global/Stepper.hh +++ b/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_; } diff --git a/src/celeritas/random/RngReseed.cc b/src/celeritas/random/RngReseed.cc index 578bb0878a..e6c5910259 100644 --- a/src/celeritas/random/RngReseed.cc +++ b/src/celeritas/random/RngReseed.cc @@ -24,14 +24,18 @@ namespace celeritas */ void reseed_rng(HostCRef const& params, HostRef 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(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; } } diff --git a/src/celeritas/random/RngReseed.cu b/src/celeritas/random/RngReseed.cu index d8675b2b1a..267492ab24 100644 --- a/src/celeritas/random/RngReseed.cu +++ b/src/celeritas/random/RngReseed.cu @@ -26,7 +26,7 @@ namespace */ __global__ void reseed_rng_kernel(DeviceCRef const params, DeviceRef 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 const params, */ void reseed_rng(DeviceCRef const& params, DeviceRef 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()); } //---------------------------------------------------------------------------// diff --git a/src/celeritas/random/RngReseed.hh b/src/celeritas/random/RngReseed.hh index b41a6aff0f..f07154dfc1 100644 --- a/src/celeritas/random/RngReseed.hh +++ b/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 const&, DeviceRef const&, - size_type); + UniqueEventId); void reseed_rng(HostCRef const&, HostRef const&, - size_type); + UniqueEventId); #if !CELER_USE_DEVICE //---------------------------------------------------------------------------// @@ -33,7 +34,7 @@ void reseed_rng(HostCRef const&, */ inline void reseed_rng(DeviceCRef const&, DeviceRef const&, - size_type) + UniqueEventId) { CELER_ASSERT_UNREACHABLE(); } diff --git a/test/celeritas/random/RngReseed.test.cc b/test/celeritas/random/RngReseed.test.cc index a9bd9a3ef7..10d1b83f88 100644 --- a/test/celeritas/random/RngReseed.test.cc +++ b/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 values; for (auto i : range(states.size()).step(128u)) {