From 1e4aa732ee57f33f5c9d716c32e9d857513d3c49 Mon Sep 17 00:00:00 2001 From: Amanda Lund Date: Fri, 16 Aug 2024 07:49:23 -0600 Subject: [PATCH] Properly clear state counters and data for reuse (#1367) * Add a helper method to reset state data * Add a generic function to fill a collection with a sequence * Move track shuffle util to global/detail --- app/celer-sim/Transporter.cc | 8 ++++ src/celeritas/CMakeLists.txt | 1 + src/celeritas/global/CoreState.cc | 25 +++++++++-- src/celeritas/global/CoreState.hh | 3 ++ src/celeritas/global/CoreTrackData.cc | 12 +++--- src/celeritas/global/Stepper.hh | 3 ++ src/celeritas/global/detail/TrackSlotUtils.cc | 36 ++++++++++++++++ src/celeritas/global/detail/TrackSlotUtils.cu | 42 +++++++++++++++++++ src/celeritas/global/detail/TrackSlotUtils.hh | 42 +++++++++++++++++++ src/celeritas/optical/TrackData.cc | 2 +- src/celeritas/track/SimData.hh | 6 --- src/celeritas/track/TrackInitData.hh | 14 ++----- src/celeritas/track/detail/TrackSortUtils.cc | 27 ------------ src/celeritas/track/detail/TrackSortUtils.cu | 40 ------------------ src/celeritas/track/detail/TrackSortUtils.hh | 39 ----------------- src/corecel/data/CollectionAlgorithms.hh | 21 ++++++++++ test/celeritas/track/Sim.test.cc | 1 + test/corecel/data/Collection.test.cc | 21 ++++++++-- 18 files changed, 208 insertions(+), 135 deletions(-) create mode 100644 src/celeritas/global/detail/TrackSlotUtils.cc create mode 100644 src/celeritas/global/detail/TrackSlotUtils.cu create mode 100644 src/celeritas/global/detail/TrackSlotUtils.hh diff --git a/app/celer-sim/Transporter.cc b/app/celer-sim/Transporter.cc index 0b00bba909..cd63096f5e 100644 --- a/app/celer-sim/Transporter.cc +++ b/app/celer-sim/Transporter.cc @@ -160,6 +160,14 @@ auto Transporter::operator()(SpanConstPrimary primaries) -> TransporterResult result.num_aborted = track_counts.alive + track_counts.queued; result.num_track_slots = stepper_->state().size(); + + if (result.num_aborted > 0) + { + // Reset the state data for the next event if the stepping loop was + // aborted early + step.reset_state(); + } + return result; } diff --git a/src/celeritas/CMakeLists.txt b/src/celeritas/CMakeLists.txt index 5cf1f1af52..e49f94ba5f 100644 --- a/src/celeritas/CMakeLists.txt +++ b/src/celeritas/CMakeLists.txt @@ -261,6 +261,7 @@ celeritas_polysource(global/alongstep/AlongStepGeneralLinearAction) celeritas_polysource(global/alongstep/AlongStepNeutralAction) celeritas_polysource(global/alongstep/AlongStepUniformMscAction) celeritas_polysource(global/alongstep/AlongStepRZMapFieldMscAction) +celeritas_polysource(global/detail/TrackSlotUtils) celeritas_polysource(neutron/model/ChipsNeutronElasticModel) celeritas_polysource(neutron/model/NeutronInelasticModel) celeritas_polysource(optical/detail/CerenkovOffloadAction) diff --git a/src/celeritas/global/CoreState.cc b/src/celeritas/global/CoreState.cc index 57af0b6d1d..726a6b7c9c 100644 --- a/src/celeritas/global/CoreState.cc +++ b/src/celeritas/global/CoreState.cc @@ -7,11 +7,11 @@ //---------------------------------------------------------------------------// #include "CoreState.hh" +#include "corecel/data/CollectionAlgorithms.hh" #include "corecel/data/Copier.hh" #include "corecel/io/Logger.hh" #include "corecel/sys/ScopedProfiling.hh" #include "celeritas/track/TrackInitParams.hh" -#include "celeritas/track/detail/TrackSortUtils.hh" #include "ActionRegistry.hh" #include "CoreParams.hh" @@ -43,8 +43,6 @@ CoreState::CoreState(CoreParams const& params, params.host_ref(), stream_id, num_track_slots); counters_.num_vacancies = num_track_slots; - counters_.num_primaries = 0; - counters_.num_initializers = 0; if constexpr (M == MemSpace::device) { @@ -110,6 +108,27 @@ Range CoreState::get_action_range(ActionId action_id) const return {thread_offsets[action_id], thread_offsets[action_id + 1]}; } +//---------------------------------------------------------------------------// +/*! + * Reset the state data. + * + * This clears the state counters and initializes the necessary state data so + * the state can be reused for a new event. This should only be necessary if + * the previous event aborted early. + */ +template +void CoreState::reset() +{ + counters_ = CoreStateCounters{}; + counters_.num_vacancies = this->size(); + + // Reset all the track slots to inactive + fill(TrackStatus::inactive, &this->ref().sim.status); + + // Mark all the track slots as empty + fill_sequence(&this->ref().init.vacancies, this->stream_id()); +} + //---------------------------------------------------------------------------// // EXPLICIT INSTANTIATION //---------------------------------------------------------------------------// diff --git a/src/celeritas/global/CoreState.hh b/src/celeritas/global/CoreState.hh index d6bbdc2dac..ce9f86bbeb 100644 --- a/src/celeritas/global/CoreState.hh +++ b/src/celeritas/global/CoreState.hh @@ -115,6 +115,9 @@ class CoreState final : public CoreStateInterface //! Get a native-memspace pointer to the mutable state data Ptr ptr() { return ptr_; } + //! Reset the state data + void reset(); + //// COUNTERS //// //! Track initialization counters diff --git a/src/celeritas/global/CoreTrackData.cc b/src/celeritas/global/CoreTrackData.cc index 0e907fc6e0..2b98077b92 100644 --- a/src/celeritas/global/CoreTrackData.cc +++ b/src/celeritas/global/CoreTrackData.cc @@ -7,7 +7,9 @@ //---------------------------------------------------------------------------// #include "CoreTrackData.hh" -#include "celeritas/track/detail/TrackSortUtils.hh" +#include "corecel/data/CollectionAlgorithms.hh" + +#include "detail/TrackSlotUtils.hh" namespace celeritas { @@ -38,18 +40,16 @@ void resize(CoreStateData* state, resize(&state->physics, params.physics, size); resize(&state->rng, params.rng, stream_id, size); resize(&state->sim, size); - resize(&state->init, params.init, size); + resize(&state->init, params.init, stream_id, size); state->stream_id = stream_id; if (params.init.track_order != TrackOrder::unsorted) { resize(&state->track_slots, size); - Span track_slots{ - state->track_slots[AllItems{}]}; - detail::fill_track_slots(track_slots, stream_id); + fill_sequence(&state->track_slots, stream_id); if (params.init.track_order == TrackOrder::shuffled) { - detail::shuffle_track_slots(track_slots, stream_id); + detail::shuffle_track_slots(&state->track_slots, stream_id); } } diff --git a/src/celeritas/global/Stepper.hh b/src/celeritas/global/Stepper.hh index af4d467981..22b974b4ae 100644 --- a/src/celeritas/global/Stepper.hh +++ b/src/celeritas/global/Stepper.hh @@ -155,6 +155,9 @@ class Stepper final : public StepperInterface //! Get the core state interface for diagnostic output CoreStateInterface const& state() const final { return state_; } + //! Reset the core state counters and data so it can be reused + void reset_state() { state_.reset(); } + private: // Params data std::shared_ptr params_; diff --git a/src/celeritas/global/detail/TrackSlotUtils.cc b/src/celeritas/global/detail/TrackSlotUtils.cc new file mode 100644 index 0000000000..b901db7d37 --- /dev/null +++ b/src/celeritas/global/detail/TrackSlotUtils.cc @@ -0,0 +1,36 @@ +//----------------------------------*-C++-*----------------------------------// +// Copyright 2024 UT-Battelle, LLC, and other Celeritas developers. +// See the top-level COPYRIGHT file for details. +// SPDX-License-Identifier: (Apache-2.0 OR MIT) +//---------------------------------------------------------------------------// +//! \file celeritas/global/detail/TrackSlotUtils.cc +//---------------------------------------------------------------------------// +#include "TrackSlotUtils.hh" + +#include +#include + +namespace celeritas +{ +namespace detail +{ +//---------------------------------------------------------------------------// +/*! + * Shuffle track slot indices. + */ +void shuffle_track_slots( + Collection* + track_slots, + StreamId) +{ + CELER_EXPECT(track_slots); + auto* start + = static_cast(track_slots->data().get()); + auto seed = static_cast(track_slots->size()); + std::mt19937 g{seed}; + std::shuffle(start, start + track_slots->size(), g); +} + +//---------------------------------------------------------------------------// +} // namespace detail +} // namespace celeritas diff --git a/src/celeritas/global/detail/TrackSlotUtils.cu b/src/celeritas/global/detail/TrackSlotUtils.cu new file mode 100644 index 0000000000..476d5c5c67 --- /dev/null +++ b/src/celeritas/global/detail/TrackSlotUtils.cu @@ -0,0 +1,42 @@ +//---------------------------------*-CUDA-*----------------------------------// +// Copyright 2024 UT-Battelle, LLC, and other Celeritas developers. +// See the top-level COPYRIGHT file for details. +// SPDX-License-Identifier: (Apache-2.0 OR MIT) +//---------------------------------------------------------------------------// +//! \file celeritas/global/detail/TrackSlotUtils.cu +//---------------------------------------------------------------------------// +#include "TrackSlotUtils.hh" + +#include +#include +#include +#include + +#include "corecel/sys/Thrust.device.hh" + +namespace celeritas +{ +namespace detail +{ +//---------------------------------------------------------------------------// +/*! + * Shuffle track slot indices. + */ +void shuffle_track_slots( + Collection* + track_slots, + StreamId stream) +{ + CELER_EXPECT(track_slots); + using result_type = thrust::default_random_engine::result_type; + thrust::default_random_engine g{ + static_cast(track_slots->size())}; + auto start = thrust::device_pointer_cast(track_slots->data().get()); + thrust::shuffle( + thrust_execute_on(stream), start, start + track_slots->size(), g); + CELER_DEVICE_CHECK_ERROR(); +} + +//---------------------------------------------------------------------------// +} // namespace detail +} // namespace celeritas diff --git a/src/celeritas/global/detail/TrackSlotUtils.hh b/src/celeritas/global/detail/TrackSlotUtils.hh new file mode 100644 index 0000000000..41423c8996 --- /dev/null +++ b/src/celeritas/global/detail/TrackSlotUtils.hh @@ -0,0 +1,42 @@ +//----------------------------------*-C++-*----------------------------------// +// Copyright 2024 UT-Battelle, LLC, and other Celeritas developers. +// See the top-level COPYRIGHT file for details. +// SPDX-License-Identifier: (Apache-2.0 OR MIT) +//---------------------------------------------------------------------------// +//! \file celeritas/global/detail/TrackSlotUtils.hh +//---------------------------------------------------------------------------// +#pragma once + +#include "corecel/Assert.hh" +#include "corecel/Macros.hh" +#include "corecel/Types.hh" +#include "corecel/data/Collection.hh" +#include "corecel/sys/ThreadId.hh" + +namespace celeritas +{ +namespace detail +{ +//---------------------------------------------------------------------------// +// Shuffle track slot indices +void shuffle_track_slots( + Collection*, + StreamId); +void shuffle_track_slots( + Collection*, + StreamId); + +//---------------------------------------------------------------------------// +// INLINE DEFINITIONS +//---------------------------------------------------------------------------// +#if !CELER_USE_DEVICE +inline void shuffle_track_slots( + Collection*, + StreamId) +{ + CELER_NOT_CONFIGURED("CUDA or HIP"); +} +#endif +//---------------------------------------------------------------------------// +} // namespace detail +} // namespace celeritas diff --git a/src/celeritas/optical/TrackData.cc b/src/celeritas/optical/TrackData.cc index e02c01bc78..bb45b02f44 100644 --- a/src/celeritas/optical/TrackData.cc +++ b/src/celeritas/optical/TrackData.cc @@ -39,7 +39,7 @@ void resize(CoreStateData* state, resize(&state->physics, params.physics, size); resize(&state->rng, params.rng, stream_id, size); resize(&state->sim, size); - resize(&state->init, params.init, size); + resize(&state->init, params.init, stream_id, size); state->stream_id = stream_id; CELER_ENSURE(state); diff --git a/src/celeritas/track/SimData.hh b/src/celeritas/track/SimData.hh index d774c86021..3c6363647c 100644 --- a/src/celeritas/track/SimData.hh +++ b/src/celeritas/track/SimData.hh @@ -183,15 +183,9 @@ void resize(SimStateData* data, size_type size) resize(&data->track_ids, size); resize(&data->parent_ids, size); resize(&data->event_ids, size); - resize(&data->num_steps, size); - fill(size_type{0}, &data->num_steps); - resize(&data->num_looping_steps, size); - fill(size_type{0}, &data->num_looping_steps); - resize(&data->time, size); - fill(real_type{0}, &data->time); resize(&data->status, size); fill(TrackStatus::inactive, &data->status); diff --git a/src/celeritas/track/TrackInitData.hh b/src/celeritas/track/TrackInitData.hh index bcef529894..b0f10588d6 100644 --- a/src/celeritas/track/TrackInitData.hh +++ b/src/celeritas/track/TrackInitData.hh @@ -10,6 +10,7 @@ #include "corecel/Types.hh" #include "corecel/cont/Range.hh" #include "corecel/data/Collection.hh" +#include "corecel/data/CollectionAlgorithms.hh" #include "corecel/data/CollectionBuilder.hh" #include "corecel/sys/Device.hh" #include "corecel/sys/ThreadId.hh" @@ -145,9 +146,6 @@ struct TrackInitStateData } }; -using TrackInitStateDeviceRef = DeviceRef; -using TrackInitStateHostRef = HostRef; - //---------------------------------------------------------------------------// /*! * Resize and initialize track initializer data. @@ -162,6 +160,7 @@ using TrackInitStateHostRef = HostRef; template void resize(TrackInitStateData* data, HostCRef const& params, + StreamId stream, size_type size) { CELER_EXPECT(params); @@ -177,13 +176,8 @@ void resize(TrackInitStateData* data, fill(size_type(0), &data->track_counters); // Initialize vacancies to mark all track slots as empty - StateCollection vacancies; - resize(&vacancies, size); - for (auto i : range(size)) - { - vacancies[TrackSlotId{i}] = TrackSlotId{i}; - } - data->vacancies = std::move(vacancies); + resize(&data->vacancies, size); + fill_sequence(&data->vacancies, stream); // Reserve space for initializers resize(&data->initializers, params.capacity); diff --git a/src/celeritas/track/detail/TrackSortUtils.cc b/src/celeritas/track/detail/TrackSortUtils.cc index c232e5de41..dfd43cb4cb 100644 --- a/src/celeritas/track/detail/TrackSortUtils.cc +++ b/src/celeritas/track/detail/TrackSortUtils.cc @@ -10,7 +10,6 @@ #include #include #include -#include #include "corecel/data/Collection.hh" @@ -65,32 +64,6 @@ IdLess(ObserverPtr) -> IdLess; //---------------------------------------------------------------------------// } // namespace -//---------------------------------------------------------------------------// -/*! - * Initialize default threads to track_slots mapping. - * - * This sets \code track_slots[i] = i \endcode . - */ -template<> -void fill_track_slots(Span track_slots, - StreamId) -{ - std::iota(track_slots.data(), track_slots.data() + track_slots.size(), 0); -} - -//---------------------------------------------------------------------------// -/*! - * Shuffle track slots. - */ -template<> -void shuffle_track_slots( - Span track_slots, StreamId) -{ - auto seed = static_cast(track_slots.size()); - std::mt19937 g{seed}; - std::shuffle(track_slots.begin(), track_slots.end(), g); -} - //---------------------------------------------------------------------------// /*! * Sort or partition tracks. diff --git a/src/celeritas/track/detail/TrackSortUtils.cu b/src/celeritas/track/detail/TrackSortUtils.cu index 59a5a35ef2..22e71d42c3 100644 --- a/src/celeritas/track/detail/TrackSortUtils.cu +++ b/src/celeritas/track/detail/TrackSortUtils.cu @@ -11,9 +11,6 @@ #include #include #include -#include -#include -#include #include #include "corecel/Macros.hh" @@ -141,43 +138,6 @@ tracks_per_action_kernel(ObserverPtr actions, //---------------------------------------------------------------------------// } // namespace -//---------------------------------------------------------------------------// -/*! - * Initialize default threads to track_slots mapping, track_slots[i] = i. - * - * TODO: move to global/detail - */ -template<> -void fill_track_slots(Span track_slots, - StreamId stream_id) -{ - thrust::sequence( - thrust_execute_on(stream_id), - thrust::device_pointer_cast(track_slots.data()), - thrust::device_pointer_cast(track_slots.data() + track_slots.size()), - 0); - CELER_DEVICE_CHECK_ERROR(); -} - -//---------------------------------------------------------------------------// -/*! - * Shuffle track slots. - * - * TODO: move to global/detail - */ -template<> -void shuffle_track_slots( - Span track_slots, StreamId stream_id) -{ - using result_type = thrust::default_random_engine::result_type; - thrust::default_random_engine g{ - static_cast(track_slots.size())}; - auto start = thrust::device_pointer_cast(track_slots.data()); - thrust::shuffle( - thrust_execute_on(stream_id), start, start + track_slots.size(), g); - CELER_DEVICE_CHECK_ERROR(); -} - //---------------------------------------------------------------------------// /*! * Sort or partition tracks. diff --git a/src/celeritas/track/detail/TrackSortUtils.hh b/src/celeritas/track/detail/TrackSortUtils.hh index f84941718b..4c4dbda9ca 100644 --- a/src/celeritas/track/detail/TrackSortUtils.hh +++ b/src/celeritas/track/detail/TrackSortUtils.hh @@ -25,31 +25,6 @@ namespace detail // HELPER FUNCTIONS //---------------------------------------------------------------------------// -// Initialize default threads to track_slots mapping, track_slots[i] = i -// TODO: move to global/detail and overload using ObserverPtr -template -void fill_track_slots(Span track_slots, StreamId); - -template<> -void fill_track_slots(Span track_slots, - StreamId); -template<> -void fill_track_slots(Span track_slots, - StreamId); - -//---------------------------------------------------------------------------// -// Shuffle tracks -// TODO: move to global/detail and overload using ObserverPtr -template -void shuffle_track_slots(Span track_slots, StreamId); - -template<> -void shuffle_track_slots( - Span track_slots, StreamId); -template<> -void shuffle_track_slots( - Span track_slots, StreamId); - //---------------------------------------------------------------------------// // Sort or partition tracks void sort_tracks(HostRef const&, TrackOrder); @@ -120,20 +95,6 @@ get_action_ptr(CoreStateData const& states, TrackOrder order) // INLINE DEFINITIONS //---------------------------------------------------------------------------// #if !CELER_USE_DEVICE -template<> -inline void -fill_track_slots(Span, StreamId) -{ - CELER_NOT_CONFIGURED("CUDA or HIP"); -} - -template<> -inline void -shuffle_track_slots(Span, StreamId) -{ - CELER_NOT_CONFIGURED("CUDA or HIP"); -} - inline void sort_tracks(DeviceRef const&, TrackOrder) { CELER_NOT_CONFIGURED("CUDA or HIP"); diff --git a/src/corecel/data/CollectionAlgorithms.hh b/src/corecel/data/CollectionAlgorithms.hh index ae88642cfa..b2dfcb3d31 100644 --- a/src/corecel/data/CollectionAlgorithms.hh +++ b/src/corecel/data/CollectionAlgorithms.hh @@ -7,8 +7,12 @@ //---------------------------------------------------------------------------// #pragma once +#include +#include + #include "Collection.hh" #include "Copier.hh" + #include "detail/Filler.hh" namespace celeritas @@ -27,6 +31,23 @@ void fill(T&& value, Collection* col) fill_impl((*col)[AllItems{}]); } +//---------------------------------------------------------------------------// +/*! + * Fill the collection with sequentially increasing values starting from zero. + */ +template +void fill_sequence(Collection* dst, StreamId stream) +{ + static_assert(W != Ownership::const_reference, + "const references cannot be filled"); + CELER_EXPECT(dst); + + std::vector src(dst->size()); + std::iota(src.begin(), src.end(), T{0}); + Copier copy{(*dst)[AllItems{}], stream}; + copy(MemSpace::host, make_span(src)); +} + //---------------------------------------------------------------------------// /*! * Copy from the given collection to host. diff --git a/test/celeritas/track/Sim.test.cc b/test/celeritas/track/Sim.test.cc index 24cfe9b290..b36a5926c9 100644 --- a/test/celeritas/track/Sim.test.cc +++ b/test/celeritas/track/Sim.test.cc @@ -84,6 +84,7 @@ TEST_F(SimTest, looping) // Check looping threshold parameters for each particle SimTrackView sim(this->sim()->host_ref(), sim_state_.ref(), TrackSlotId{0}); + sim = SimTrackInitializer{TrackId{0}, TrackId{}, EventId{0}, 0}; for (auto pid : range(ParticleId{this->particle()->size()})) { auto const& looping = sim.looping_threshold(pid); diff --git a/test/corecel/data/Collection.test.cc b/test/corecel/data/Collection.test.cc index 3ee028fd1e..0d963b19dc 100644 --- a/test/corecel/data/Collection.test.cc +++ b/test/corecel/data/Collection.test.cc @@ -328,11 +328,21 @@ TEST_F(SimpleCollectionTest, algo_host) fill(123, &src); EXPECT_EQ(123, src[IntId{0}]); EXPECT_EQ(123, src[IntId{3}]); - src[IntId{1}] = 2; + + // Test 'fill_sequence' + fill_sequence(&src, StreamId{}); + for (size_type i = 0; i < src.size(); ++i) + { + EXPECT_EQ(i, src[IntId{i}]); + } // Test 'copy_to_host' std::vector dst(src.size()); copy_to_host(src, make_span(dst)); + for (size_type i = 0; i < src.size(); ++i) + { + EXPECT_EQ(i, dst[i]); + } } TEST_F(SimpleCollectionTest, TEST_IF_CELER_DEVICE(algo_device)) @@ -349,11 +359,16 @@ TEST_F(SimpleCollectionTest, TEST_IF_CELER_DEVICE(algo_device)) EXPECT_TRUE((std::is_same_v{}]), LdgSpan>)); + // Test 'fill_sequence' + fill_sequence(&src, StreamId{}); + // Test 'copy_to_host' std::vector dst(src.size()); copy_to_host(src, make_span(dst)); - EXPECT_EQ(123, dst.front()); - EXPECT_EQ(123, dst.back()); + for (size_type i = 0; i < src.size(); ++i) + { + EXPECT_EQ(i, dst[i]); + } } //---------------------------------------------------------------------------//