From 412c9ef7f807c3b7b38a5310b9b4fb2b7540f08a Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Tue, 25 Nov 2025 10:29:23 -0500 Subject: [PATCH 01/18] Reapply "Prohibit ratracing from outside" This reverts commit b4ff3bbb01eacc53ca8326d152f1f5177df73454. --- src/geocel/vg/VecgeomTrackView.hh | 39 ++++++---------------------- src/geocel/vg/detail/BVHNavigator.hh | 13 +++++++--- 2 files changed, 17 insertions(+), 35 deletions(-) diff --git a/src/geocel/vg/VecgeomTrackView.hh b/src/geocel/vg/VecgeomTrackView.hh index 682694767f..48584cf051 100644 --- a/src/geocel/vg/VecgeomTrackView.hh +++ b/src/geocel/vg/VecgeomTrackView.hh @@ -426,32 +426,12 @@ CELER_FUNCTION Real3 VecgeomTrackView::normal() const /*! * Find the distance to the next geometric boundary. * - * This function is allowed to be called from the exterior for ray tracing. + * \todo To support ray tracing from unknown distances from the world, we + * could add another special function \c find_first_step . */ CELER_FUNCTION Propagation VecgeomTrackView::find_next_step() { - if (this->is_outside()) - { - // SPECIAL CASE: find distance to interior from outside world volume - auto const& world_pv = this->world(); - next_step_ = world_pv.DistanceToIn(detail::to_vector(pos_), - detail::to_vector(dir_), - vecgeom::kInfLength); - - next_step_ = max(next_step_, this->extra_push()); - vgnext_.Clear(); - Propagation result; - result.distance = next_step_; - result.boundary = next_step_ < vecgeom::kInfLength; - - if (result.boundary) - { - vgnext_.Push(&world_pv); - vgnext_.SetBoundaryState(true); - } - - return result; - } + CELER_EXPECT(!this->is_outside()); return this->find_next_step(vecgeom::kInfLength); } @@ -571,22 +551,19 @@ CELER_FUNCTION void VecgeomTrackView::move_to_boundary() */ CELER_FUNCTION void VecgeomTrackView::cross_boundary() { + CELER_EXPECT(!this->is_outside()); CELER_EXPECT(this->is_on_boundary()); CELER_EXPECT(this->is_next_boundary()); // Relocate to next tracking volume (maybe across multiple boundaries) if (vgnext_.Top() != nullptr) { - if (!CELERITAS_VECGEOM_SURFACE || !vgstate_.IsOutside()) - { - // In surf model, relocation does not work from [OUTSIDE] - Navigator::RelocateToNextVolume(detail::to_vector(this->pos_), - detail::to_vector(this->dir_), + Navigator::RelocateToNextVolume(detail::to_vector(this->pos_), + detail::to_vector(this->dir_), #if CELERITAS_VECGEOM_SURFACE - *next_surf_, + *next_surf_, #endif - vgnext_); - } + vgnext_); } vgstate_ = vgnext_; diff --git a/src/geocel/vg/detail/BVHNavigator.hh b/src/geocel/vg/detail/BVHNavigator.hh index 065d9e2285..0565cd82bd 100644 --- a/src/geocel/vg/detail/BVHNavigator.hh +++ b/src/geocel/vg/detail/BVHNavigator.hh @@ -19,7 +19,9 @@ #include #include +#include "corecel/Macros.hh" #include "corecel/math/Algorithms.hh" +#include "geocel/vg/VecgeomTypes.hh" #ifdef VECGEOM_ENABLE_CUDA # include @@ -31,9 +33,6 @@ # include "VgNavStateWrapper.hh" #endif -#include "corecel/Macros.hh" -#include "geocel/vg/VecgeomTypes.hh" - namespace celeritas { namespace detail @@ -73,6 +72,7 @@ class BVHNavigator } } + CELER_ASSERT(vol); path.Push(vol); VgReal3 currentpoint(point); @@ -93,8 +93,12 @@ class BVHNavigator break; } - currentpoint = daughterlocalpoint; + CELER_ASSERT(vol); path.Push(vol); + + // Transform to daughter + currentpoint = daughterlocalpoint; + // Only exclude the placed volume once since we could enter it // again via a different volume history. exclude = nullptr; @@ -104,6 +108,7 @@ class BVHNavigator CELER_FUNCTION static void RelocatePoint(VgReal3 const& localpoint, NavState& path) { + CELER_EXPECT(!path.IsOutside()); VgPlacedVol const* currentmother = path.Top(); VgReal3 transformed = localpoint; do From bc996112e61ace688e2e9c26c5ad922f50cfef75 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Tue, 25 Nov 2025 09:49:07 -0500 Subject: [PATCH 02/18] Add safety with lmit --- src/geocel/vg/VecgeomTrackView.hh | 9 ++------- src/geocel/vg/detail/BVHNavigator.hh | 8 +++----- src/geocel/vg/detail/SurfNavigator.hh | 7 ++++--- 3 files changed, 9 insertions(+), 15 deletions(-) diff --git a/src/geocel/vg/VecgeomTrackView.hh b/src/geocel/vg/VecgeomTrackView.hh index 48584cf051..6353a095bc 100644 --- a/src/geocel/vg/VecgeomTrackView.hh +++ b/src/geocel/vg/VecgeomTrackView.hh @@ -512,13 +512,8 @@ CELER_FUNCTION real_type VecgeomTrackView::find_safety(real_type max_radius) CELER_EXPECT(!this->is_on_boundary()); CELER_EXPECT(max_radius > 0); - real_type safety = Navigator::ComputeSafety(detail::to_vector(this->pos()), - vgstate_ -#if VECGEOM_VERSION >= 0x200000 - , - max_radius -#endif - ); + real_type safety = Navigator::ComputeSafety( + detail::to_vector(this->pos()), vgstate_, max_radius); safety = min(safety, max_radius); // Since the reported "safety" is negative if we've moved slightly beyond diff --git a/src/geocel/vg/detail/BVHNavigator.hh b/src/geocel/vg/detail/BVHNavigator.hh index 0565cd82bd..ffad2c5367 100644 --- a/src/geocel/vg/detail/BVHNavigator.hh +++ b/src/geocel/vg/detail/BVHNavigator.hh @@ -240,11 +240,9 @@ class BVHNavigator public: // Computes the isotropic safety from the globalpoint. - CELER_FUNCTION static double - ComputeSafety(VgReal3 const& globalpoint, - NavState const& state, - vg_real_type safety - = std::numeric_limits::infinity()) + CELER_FUNCTION static double ComputeSafety(VgReal3 const& globalpoint, + NavState const& state, + vg_real_type safety) { VgPlacedVol const* pvol = state.Top(); vecgeom::Transformation3D m; diff --git a/src/geocel/vg/detail/SurfNavigator.hh b/src/geocel/vg/detail/SurfNavigator.hh index d1abe958bf..929664c672 100644 --- a/src/geocel/vg/detail/SurfNavigator.hh +++ b/src/geocel/vg/detail/SurfNavigator.hh @@ -70,12 +70,13 @@ class SurfNavigator /// @param globalpoint Point in global coordinates /// @param state Path where to compute safety /// @return Isotropic safe distance - CELER_FUNCTION static vg_real_type - ComputeSafety(VgReal3 const& globalpoint, NavState const& state) + CELER_FUNCTION static vg_real_type ComputeSafety(VgReal3 const& globalpoint, + NavState const& state, + Precision limit) { auto safety = vgbrep::protonav::BVHSurfNavigator::ComputeSafety( - globalpoint, state); + globalpoint, state, limit); return safety; } From 272f0986c4be7c993a58958786f20d4c00cf908b Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Tue, 25 Nov 2025 11:05:20 -0500 Subject: [PATCH 03/18] Remove world accessor --- src/geocel/vg/VecgeomTrackView.hh | 20 ++++---------------- 1 file changed, 4 insertions(+), 16 deletions(-) diff --git a/src/geocel/vg/VecgeomTrackView.hh b/src/geocel/vg/VecgeomTrackView.hh index 6353a095bc..ace78faeee 100644 --- a/src/geocel/vg/VecgeomTrackView.hh +++ b/src/geocel/vg/VecgeomTrackView.hh @@ -193,9 +193,6 @@ class VecgeomTrackView // Whether the next distance-to-boundary is to a surface inline CELER_FUNCTION bool is_next_boundary() const; - // Get a reference to the world volume - inline CELER_FUNCTION VgPlacedVol const& world() const; - // Get a reference to the current volume instance inline CELER_FUNCTION VgPlacedVol const& physical_volume() const; @@ -285,9 +282,11 @@ VecgeomTrackView::operator=(Initializer_t const& init) // Set up current state and locate daughter volume vgstate_.Clear(); #if CELERITAS_VECGEOM_SURFACE - auto world = vecgeom::NavigationState::WorldId(); + // VecGeom's BVHSurfNav takes `int pvol_id ` but vecgeom's navtuple/index + // return NavIndex_t via `NavInd` :( + VgPlacedVolumeInt world = vecgeom::NavigationState::WorldId(); #else - auto const* world = &this->world(); + auto const* world = params_.scalars.world(); #endif // LocatePointIn sets `vgstate_` constexpr bool contains_point = true; @@ -645,17 +644,6 @@ CELER_FUNCTION bool VecgeomTrackView::is_next_boundary() const } } -//---------------------------------------------------------------------------// -/*! - * Get a reference to the world volume instance. - */ -CELER_FUNCTION auto VecgeomTrackView::world() const -> VgPlacedVol const& -{ - auto* pv = params_.scalars.world(); - CELER_ENSURE(pv); - return *pv; -} - //---------------------------------------------------------------------------// /*! * Get a reference to the current volume. From c84c4bc03c8189f5a16f0c9b5410f59df8ef11a2 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Tue, 25 Nov 2025 13:26:43 -0500 Subject: [PATCH 04/18] Delete unused methods --- src/geocel/vg/detail/BVHNavigator.hh | 317 ++++++-------------------- src/geocel/vg/detail/SurfNavigator.hh | 28 +-- 2 files changed, 77 insertions(+), 268 deletions(-) diff --git a/src/geocel/vg/detail/BVHNavigator.hh b/src/geocel/vg/detail/BVHNavigator.hh index ffad2c5367..600e6afbe1 100644 --- a/src/geocel/vg/detail/BVHNavigator.hh +++ b/src/geocel/vg/detail/BVHNavigator.hh @@ -105,140 +105,6 @@ class BVHNavigator } } - CELER_FUNCTION static void - RelocatePoint(VgReal3 const& localpoint, NavState& path) - { - CELER_EXPECT(!path.IsOutside()); - VgPlacedVol const* currentmother = path.Top(); - VgReal3 transformed = localpoint; - do - { - path.Pop(); - transformed = currentmother->GetTransformation()->InverseTransform( - transformed); - currentmother = path.Top(); - } while (currentmother - && (currentmother->IsAssembly() - || !currentmother->UnplacedContains(transformed))); - - if (currentmother) - { - path.Pop(); - return LocatePointIn(currentmother, transformed, path, false); - } - } - - private: - // Computes a step in the current volume from the localpoint into localdir, - // taking step_limit into account. If a volume is hit, the function calls - // out_state.SetBoundaryState(true) and hitcandidate is set to the hit - // daughter volume, or kept unchanged if the current volume is left. - CELER_FUNCTION static double - ComputeStepAndHit(VgReal3 const& localpoint, - VgReal3 const& localdir, - vg_real_type step_limit, - NavState const& in_state, - NavState& out_state, - VgPlacedVol const*& hitcandidate) - { - if (step_limit <= 0) - { - // We don't need to ask any solid, step not limited by geometry. - in_state.CopyTo(&out_state); - out_state.SetBoundaryState(false); - return 0; - } - - vg_real_type step = step_limit; - VgPlacedVol const* pvol = in_state.Top(); - - // need to calc DistanceToOut first - step = pvol->DistanceToOut(localpoint, localdir, step_limit); - - if (step < 0) - step = 0; - - if (pvol->GetDaughters().size() > 0) - { - auto bvh - = vecgeom::BVHManager::GetBVH(pvol->GetLogicalVolume()->id()); - bvh->CheckDaughterIntersections( - localpoint, localdir, step, pvol, hitcandidate); - } - - // now we have the candidates and we prepare the out_state - in_state.CopyTo(&out_state); - if (step == vecgeom::kInfLength && step_limit > 0) - { - out_state.SetBoundaryState(true); - do - { - out_state.Pop(); - } while (out_state.Top()->IsAssembly()); - - return vecgeom::kTolerance; - } - - // Is geometry further away than physics step? - if (step > step_limit) - { - // Then this is a phyics step and we don't need to do anything. - out_state.SetBoundaryState(false); - return step_limit; - } - - // Otherwise it is a geometry step and we push the point to the - // boundary. - out_state.SetBoundaryState(true); - - if (step < 0) - { - step = 0; - } - - return step; - } - - // Computes a step in the current volume from the localpoint into localdir, - // until the next daughter bounding box, taking step_limit into account. - CELER_FUNCTION static double ApproachNextVolume(VgReal3 const& localpoint, - VgReal3 const& localdir, - vg_real_type step_limit, - NavState const& in_state) - { - vg_real_type step = step_limit; - VgPlacedVol const* pvol = in_state.Top(); - - if (pvol->GetDaughters().size() > 0) - { - auto bvh - = vecgeom::BVHManager::GetBVH(pvol->GetLogicalVolume()->id()); - // bvh->CheckDaughterIntersections(localpoint, localdir, step, - // pvol, hitcandidate); - bvh->ApproachNextDaughter(localpoint, localdir, step, pvol); - // Make sure we don't "step" on next boundary - step -= 10 * vecgeom::kTolerance; - } - - if (step == vecgeom::kInfLength && step_limit > 0) - return 0; - - // Is geometry further away than physics step? - if (step > step_limit) - { - // Then this is a phyics step and we don't need to do anything. - return step_limit; - } - - if (step < 0) - { - step = 0; - } - - return step; - } - - public: // Computes the isotropic safety from the globalpoint. CELER_FUNCTION static double ComputeSafety(VgReal3 const& globalpoint, NavState const& state, @@ -262,88 +128,6 @@ class BVHNavigator return safety; } - // Computes a step from the globalpoint (which must be in the current - // volume) into globaldir, taking step_limit into account. If a volume is - // hit, the function calls out_state.SetBoundaryState(true) and relocates - // the state to the next volume. - CELER_FUNCTION static double - ComputeStepAndPropagatedState(VgReal3 const& globalpoint, - VgReal3 const& globaldir, - vg_real_type step_limit, - NavState const& in_state, - NavState& out_state, - vg_real_type push = 0) - { - // If we are on the boundary, push a bit more - if (in_state.IsOnBoundary()) - { - push += kBoundaryPush; - } - if (step_limit < push) - { - // Go as far as the step limit says, assuming there is no boundary. - // TODO: Does this make sense? - in_state.CopyTo(&out_state); - out_state.SetBoundaryState(false); - return step_limit; - } - step_limit -= push; - - // calculate local point/dir from global point/dir - VgReal3 localpoint; - VgReal3 localdir; - // Impl::DoGlobalToLocalTransformation(in_state, globalpoint, - // globaldir, localpoint, localdir); - vecgeom::Transformation3D m; - in_state.TopMatrix(m); - localpoint = m.Transform(globalpoint); - localdir = m.TransformDirection(globaldir); - // The user may want to move point from boundary before computing the - // step - localpoint += push * localdir; - - VgPlacedVol const* hitcandidate = nullptr; - vg_real_type step = ComputeStepAndHit( - localpoint, localdir, step_limit, in_state, out_state, hitcandidate); - step += push; - - if (out_state.IsOnBoundary()) - { - // Relocate the point after the step to refine out_state. - localpoint += (step + kBoundaryPush) * localdir; - - if (!hitcandidate) - { - // We didn't hit a daughter but instead we're exiting the - // current volume. - RelocatePoint(localpoint, out_state); - } - else - { - // Otherwise check if we're directly entering other daughters - // transitively. - localpoint - = hitcandidate->GetTransformation()->Transform(localpoint); - LocatePointIn(hitcandidate, localpoint, out_state, false); - } - - if (out_state.Top() != nullptr) - { - while (out_state.Top()->IsAssembly() - || out_state.HasSamePathAsOther(in_state)) - { - out_state.Pop(); - } - CELER_ASSERT(!out_state.Top() - ->GetLogicalVolume() - ->GetUnplacedVolume() - ->IsAssembly()); - } - } - - return step; - } - // Computes a step from the globalpoint (which must be in the current // volume) into globaldir, taking step_limit into account. If a volume is // hit, the function calls out_state.SetBoundaryState(true) and @@ -356,9 +140,10 @@ class BVHNavigator VgReal3 const& globaldir, vg_real_type step_limit, NavState const& in_state, - NavState& out_state, - vg_real_type push = 0) + NavState& out_state) { + vg_real_type push = 0; + // If we are on the boundary, push a bit more if (in_state.IsOnBoundary()) { @@ -422,30 +207,6 @@ class BVHNavigator return step; } - // Computes a step from the globalpoint (which must be in the current - // volume) into globaldir, taking step_limit into account. - CELER_FUNCTION static vg_real_type - ComputeStepToApproachNextVolume(VgReal3 const& globalpoint, - VgReal3 const& globaldir, - vg_real_type step_limit, - NavState const& in_state) - { - // calculate local point/dir from global point/dir - VgReal3 localpoint; - VgReal3 localdir; - // Impl::DoGlobalToLocalTransformation(in_state, globalpoint, - // globaldir, localpoint, localdir); - vecgeom::Transformation3D m; - in_state.TopMatrix(m); - localpoint = m.Transform(globalpoint); - localdir = m.TransformDirection(globaldir); - - vg_real_type step - = ApproachNextVolume(localpoint, localdir, step_limit, in_state); - - return step; - } - // Relocate a state that was returned from ComputeStepAndNextVolume: It // recursively locates the pushed point in the containing volume. CELER_FUNCTION static void RelocateToNextVolume(VgReal3 const& globalpoint, @@ -477,6 +238,78 @@ class BVHNavigator ->IsAssembly()); } } + + private: + // USED ONLY BY ComputeStepAndNextVolume + // Computes a step in the current volume from the localpoint into localdir, + // taking step_limit into account. If a volume is hit, the function calls + // out_state.SetBoundaryState(true) and hitcandidate is set to the hit + // daughter volume, or kept unchanged if the current volume is left. + CELER_FUNCTION static double + ComputeStepAndHit(VgReal3 const& localpoint, + VgReal3 const& localdir, + vg_real_type step_limit, + NavState const& in_state, + NavState& out_state, + VgPlacedVol const*& hitcandidate) + { + if (step_limit <= 0) + { + // We don't need to ask any solid, step not limited by geometry. + in_state.CopyTo(&out_state); + out_state.SetBoundaryState(false); + return 0; + } + + vg_real_type step = step_limit; + VgPlacedVol const* pvol = in_state.Top(); + + // need to calc DistanceToOut first + step = pvol->DistanceToOut(localpoint, localdir, step_limit); + + if (step < 0) + step = 0; + + if (pvol->GetDaughters().size() > 0) + { + auto bvh + = vecgeom::BVHManager::GetBVH(pvol->GetLogicalVolume()->id()); + bvh->CheckDaughterIntersections( + localpoint, localdir, step, pvol, hitcandidate); + } + + // now we have the candidates and we prepare the out_state + in_state.CopyTo(&out_state); + if (step == vecgeom::kInfLength && step_limit > 0) + { + out_state.SetBoundaryState(true); + do + { + out_state.Pop(); + } while (out_state.Top()->IsAssembly()); + + return vecgeom::kTolerance; + } + + // Is geometry further away than physics step? + if (step > step_limit) + { + // Then this is a phyics step and we don't need to do anything. + out_state.SetBoundaryState(false); + return step_limit; + } + + // Otherwise it is a geometry step and we push the point to the + // boundary. + out_state.SetBoundaryState(true); + + if (step < 0) + { + step = 0; + } + + return step; + } }; //---------------------------------------------------------------------------// diff --git a/src/geocel/vg/detail/SurfNavigator.hh b/src/geocel/vg/detail/SurfNavigator.hh index 929664c672..6820ac2898 100644 --- a/src/geocel/vg/detail/SurfNavigator.hh +++ b/src/geocel/vg/detail/SurfNavigator.hh @@ -45,8 +45,6 @@ class SurfNavigator using SurfData = vgbrep::SurfData; using NavState = detail::VgNavStateWrapper; - static constexpr vg_real_type kBoundaryPush = 10 * vecgeom::kTolerance; - /// @brief Locates the point in the geometry volume tree /// @param pvol_id Placed volume id to be checked first /// @param point Point to be checked, in the local frame of pvol @@ -82,11 +80,7 @@ class SurfNavigator // Computes a step from the globalpoint (which must be in the current // volume) into globaldir, taking step_limit into account. If a volume is - // hit, the function calls out_state.SetBoundaryState(true) and relocates - // the state to the next volume. - // - // The surface model does automatic relocation, so this function does it as - // well. + // hit, the function calls out_state.SetBoundaryState(true) CELER_FUNCTION static vg_real_type ComputeStepAndNextVolume(VgReal3 const& globalpoint, VgReal3 const& globaldir, @@ -113,25 +107,7 @@ class SurfNavigator return step; } - // Computes a step from the globalpoint (which must be in the current - // volume) into globaldir, taking step_limit into account. If a volume is - // hit, the function calls out_state.SetBoundaryState(true) and relocates - // the state to the next volume. - CELER_FUNCTION static vg_real_type - ComputeStepAndPropagatedState(VgReal3 const& globalpoint, - VgReal3 const& globaldir, - vg_real_type step_limit, - VgSurfaceInt& hit_surf, - NavState const& in_state, - NavState& out_state) - { - return ComputeStepAndNextVolume( - globalpoint, globaldir, step_limit, in_state, out_state, hit_surf); - } - - // Relocate a state that was returned from ComputeStepAndNextVolume: the - // surface model does this computation within ComputeStepAndNextVolume, so - // the relocation does nothing + // Relocate a state that was returned from ComputeStepAndNextVolume CELER_FUNCTION static void RelocateToNextVolume(VgReal3 const& globalpoint, VgReal3 const& globaldir, VgSurfaceInt hitsurf_index, From 47224a03f56da9533824997381ea7c7809689424 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Fri, 6 Feb 2026 14:29:37 -0500 Subject: [PATCH 05/18] IWYU --- src/geocel/vg/VecgeomData.hh | 1 - 1 file changed, 1 deletion(-) diff --git a/src/geocel/vg/VecgeomData.hh b/src/geocel/vg/VecgeomData.hh index 4c6aadadf9..fd1e15fdd7 100644 --- a/src/geocel/vg/VecgeomData.hh +++ b/src/geocel/vg/VecgeomData.hh @@ -13,7 +13,6 @@ #include "corecel/Macros.hh" #include "corecel/Types.hh" #include "corecel/data/Collection.hh" -#include "corecel/data/CollectionBuilder.hh" #include "corecel/sys/ThreadId.hh" #include "geocel/Types.hh" From 3552ee038925aa90b3231d256dd6d60bcb980400 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Thu, 18 Dec 2025 12:31:46 -0500 Subject: [PATCH 06/18] Validate pointers from runtime in BVHNavigator --- src/geocel/vg/VecgeomParams.cc | 1 + src/geocel/vg/detail/VecgeomSetup.cu | 59 +++++++++++++++++++++++++++- src/geocel/vg/detail/VecgeomSetup.hh | 13 +++++- 3 files changed, 70 insertions(+), 3 deletions(-) diff --git a/src/geocel/vg/VecgeomParams.cc b/src/geocel/vg/VecgeomParams.cc index cc4924a105..072dd0a1b7 100644 --- a/src/geocel/vg/VecgeomParams.cc +++ b/src/geocel/vg/VecgeomParams.cc @@ -837,6 +837,7 @@ void VecgeomParams::build_volume_tracking() check_bvh_device_pointers(); check_navindex_device_pointers(); + detail::check_other_device_pointers(); device_ownership_ = Ownership::value; } diff --git a/src/geocel/vg/detail/VecgeomSetup.cu b/src/geocel/vg/detail/VecgeomSetup.cu index 600019b8dc..471cb1e713 100644 --- a/src/geocel/vg/detail/VecgeomSetup.cu +++ b/src/geocel/vg/detail/VecgeomSetup.cu @@ -6,8 +6,13 @@ //---------------------------------------------------------------------------// #include "VecgeomSetup.hh" +#include #include +#if VECGEOM_VERSION >= 0x020000 +# include +#endif + #include "corecel/data/DeviceVector.hh" #if CELERITAS_VECGEOM_SURFACE @@ -39,7 +44,7 @@ struct BvhGetter pointer_type* dest{nullptr}; - CELER_FUNCTION void operator()(ThreadId tid) + __device__ void operator()(ThreadId tid) { CELER_EXPECT(tid == ThreadId{0}); *dest = vecgeom::cuda::BVHManager::GetBVH(0); @@ -62,6 +67,46 @@ struct NavIndexGetter } }; +//---------------------------------------------------------------------------// +//! Copy the logical volume pointer table +struct LogicalVolumesGetter +{ + using pointer_type = vecgeom::cuda::LogicalVolume const*; + static constexpr char const label[] = "logical-volumes"; + + pointer_type* dest{nullptr}; + + __device__ void operator()(ThreadId tid) + { + CELER_EXPECT(tid == ThreadId{0}); +#if VECGEOM_VERSION >= 0x020000 + *dest = vecgeom::globaldevicegeomdata::gDeviceLogicalVolumes; +#else + *dest = nullptr; +#endif + } +}; + +//---------------------------------------------------------------------------// +//! Copy the logical volume pointer table +struct PlacedVolumesGetter +{ + using pointer_type = vecgeom::cuda::VPlacedVolume const*; + static constexpr char const label[] = "placed-volumes"; + + pointer_type* dest{nullptr}; + + __device__ void operator()(ThreadId tid) + { + CELER_EXPECT(tid == ThreadId{0}); +#if VECGEOM_VERSION >= 0x020000 + *dest = vecgeom::globaldevicegeomdata::gCompactPlacedVolBuffer; +#else + *dest = nullptr; +#endif + } +}; + //---------------------------------------------------------------------------// //! Launch a kernel to copy a value from global memory template @@ -144,6 +189,18 @@ CudaPointers navindex_pointers_device() return result; } +void check_other_device_pointers() +{ + if constexpr (VECGEOM_VERSION < 0x020000) + return; + + CELER_VALIDATE(get_device_pointer() != nullptr, + << "failed to copy VG logical volumes to GPU"); + + CELER_VALIDATE(get_device_pointer() != nullptr, + << "failed to copy VG places volumes to GPU"); +} + #if CELER_VGNAV == CELER_VGNAV_TUPLE //---------------------------------------------------------------------------// /* diff --git a/src/geocel/vg/detail/VecgeomSetup.hh b/src/geocel/vg/detail/VecgeomSetup.hh index 60aed02bd2..cb278b3310 100644 --- a/src/geocel/vg/detail/VecgeomSetup.hh +++ b/src/geocel/vg/detail/VecgeomSetup.hh @@ -48,7 +48,11 @@ CudaPointers bvh_pointers_device(); //---------------------------------------------------------------------------// // Get pointers to the global nav index after setup, for consistency checking -CudaPointers navindex_pointers_device(); +CudaPointers navindex_pointers_device(); + +//---------------------------------------------------------------------------// +// Check consistency of BVH pointers for device data in vecgeom 2.x +void check_other_device_pointers(); //---------------------------------------------------------------------------// // Default-initialize navigation state because DeviceVector doesn't @@ -72,7 +76,12 @@ inline CudaPointers bvh_pointers_device() CELER_ASSERT_UNREACHABLE(); } -inline CudaPointers navindex_pointers_device() +inline CudaPointers navindex_pointers_device() +{ + CELER_ASSERT_UNREACHABLE(); +} + +inline void check_other_device_pointers() { CELER_ASSERT_UNREACHABLE(); } From f9e218353c3e42deeec8cd2b22f2fcf53ee5b6c5 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Fri, 6 Feb 2026 14:36:18 -0500 Subject: [PATCH 07/18] Simplify push code --- src/geocel/vg/detail/BVHNavigator.hh | 14 ++++---------- 1 file changed, 4 insertions(+), 10 deletions(-) diff --git a/src/geocel/vg/detail/BVHNavigator.hh b/src/geocel/vg/detail/BVHNavigator.hh index 600e6afbe1..41db3e9706 100644 --- a/src/geocel/vg/detail/BVHNavigator.hh +++ b/src/geocel/vg/detail/BVHNavigator.hh @@ -142,19 +142,12 @@ class BVHNavigator NavState const& in_state, NavState& out_state) { - vg_real_type push = 0; - // If we are on the boundary, push a bit more - if (in_state.IsOnBoundary()) - { - push += kBoundaryPush; - } + vg_real_type push = in_state.IsOnBoundary() ? kBoundaryPush : 0; + if (step_limit < push) { - // Go as far as the step limit says, assuming there is no boundary. - // TODO: Does this make sense? - in_state.CopyTo(&out_state); - out_state.SetBoundaryState(false); + // Ignore small steps on boundary without a change in state return step_limit; } step_limit -= push; @@ -181,6 +174,7 @@ class BVHNavigator { if (!hitcandidate) { + // Pop back up to the parent volume VgPlacedVol const* currentmother = out_state.Top(); VgReal3 transformed = localpoint; // Push the point inside the next volume. From c4afb976ea3add0bc119ccbda2d0154d5fb282f0 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Fri, 6 Feb 2026 14:36:27 -0500 Subject: [PATCH 08/18] Add precision assertions --- src/geocel/vg/VecgeomTypes.hh | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/src/geocel/vg/VecgeomTypes.hh b/src/geocel/vg/VecgeomTypes.hh index 405a618c31..bdf1c8bfaf 100644 --- a/src/geocel/vg/VecgeomTypes.hh +++ b/src/geocel/vg/VecgeomTypes.hh @@ -38,6 +38,12 @@ # define CELER_VGNAV CELER_VGNAV_PATH #endif +#ifdef VECGEOM_SINGLE_PRECISION +static_assert(CELERITAS_REAL_TYPE == CELERITAS_REAL_TYPE_FLOAT); +#else +static_assert(CELERITAS_REAL_TYPE == CELERITAS_REAL_TYPE_DOUBLE); +#endif + namespace vecgeom { #if VECGEOM_VERSION >= 0x020000 From 49bfdc0fe7c1864e4ae0e60412dfc6032959a3ca Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Sat, 29 Nov 2025 11:01:34 -0500 Subject: [PATCH 09/18] Add length units --- src/geocel/detail/LengthUnits.hh | 3 +++ 1 file changed, 3 insertions(+) diff --git a/src/geocel/detail/LengthUnits.hh b/src/geocel/detail/LengthUnits.hh index 2adbdc7f8f..20900cad0c 100644 --- a/src/geocel/detail/LengthUnits.hh +++ b/src/geocel/detail/LengthUnits.hh @@ -23,14 +23,17 @@ namespace lengthunits CELER_ICC meter{100}; CELER_ICC centimeter{1}; CELER_ICC millimeter{0.1}; +inline constexpr char const label[] = "cm"; #elif CELERITAS_UNITS == CELERITAS_UNITS_SI CELER_ICC meter{1}; CELER_ICC centimeter{0.01}; CELER_ICC millimeter{0.001}; +inline constexpr char const label[] = "m"; #elif CELERITAS_UNITS == CELERITAS_UNITS_CLHEP CELER_ICC meter{1000}; CELER_ICC centimeter{10}; CELER_ICC millimeter{1}; +inline constexpr char const label[] = "mm"; #else # error "CELERITAS_UNITS is undefined" #endif From 5056d008fbdb9d5944aa28119a3760919ce1b95c Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Fri, 6 Feb 2026 14:42:00 -0500 Subject: [PATCH 10/18] Check for failure during initialization --- src/geocel/vg/VecgeomTrackView.hh | 22 +++++++++++++++++++++- 1 file changed, 21 insertions(+), 1 deletion(-) diff --git a/src/geocel/vg/VecgeomTrackView.hh b/src/geocel/vg/VecgeomTrackView.hh index ace78faeee..b9de7c3ae3 100644 --- a/src/geocel/vg/VecgeomTrackView.hh +++ b/src/geocel/vg/VecgeomTrackView.hh @@ -41,6 +41,12 @@ # include "detail/VgNavStateWrapper.hh" #endif +#if !CELER_DEVICE_COMPILE +# include "corecel/io/Logger.hh" +# include "corecel/io/Repr.hh" +# include "geocel/detail/LengthUnits.hh" +#endif + namespace celeritas { //---------------------------------------------------------------------------// @@ -121,7 +127,7 @@ class VecgeomTrackView // Whether the track is exactly on a surface CELER_FORCEINLINE_FUNCTION bool is_on_boundary() const; //! Whether the last operation resulted in an error - CELER_FORCEINLINE_FUNCTION bool failed() const { return false; } + CELER_FORCEINLINE_FUNCTION bool failed() const { return failed_; } // Get the normal vector of the current surface inline CELER_FUNCTION Real3 normal() const; @@ -184,6 +190,7 @@ class VecgeomTrackView // Temporary data real_type next_step_{0}; + bool failed_{false}; //// HELPER FUNCTIONS //// @@ -251,6 +258,7 @@ CELER_FUNCTION VecgeomTrackView& VecgeomTrackView::operator=(Initializer_t const& init) { CELER_EXPECT(is_soft_unit_vector(init.dir)); + failed_ = false; // Initialize direction dir_ = init.dir; @@ -292,6 +300,17 @@ VecgeomTrackView::operator=(Initializer_t const& init) constexpr bool contains_point = true; Navigator::LocatePointIn( world, detail::to_vector(pos_), vgstate_, contains_point); + + if (CELER_UNLIKELY(vgstate_.IsOutside())) + { +#if !CELER_DEVICE_COMPILE + auto msg = CELER_LOG_LOCAL(error); + msg << "Failed to initialize geometry state at " << repr(pos_) + << lengthunits::label; +#endif + failed_ = true; + } + return *this; } @@ -448,6 +467,7 @@ CELER_FUNCTION Propagation VecgeomTrackView::find_next_step(real_type max_step) { *next_surf_ = null_surface(); } + // TODO: vgnext is simply copied and the boundary flag optionally set next_step_ = Navigator::ComputeStepAndNextVolume(detail::to_vector(pos_), detail::to_vector(dir_), From a8f095975bf9de8900709e5fcb53379d80ab7491 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Fri, 6 Feb 2026 14:43:40 -0500 Subject: [PATCH 11/18] Check for next step failure --- src/geocel/vg/VecgeomTrackView.hh | 19 ++++++++++++++++++- test/celeritas/field/FieldPropagator.test.cc | 12 +++++++----- test/celeritas/geo/HeuristicGeoExecutor.hh | 7 ++++++- test/celeritas/global/Stepper.test.cc | 13 +++++++++---- 4 files changed, 40 insertions(+), 11 deletions(-) diff --git a/src/geocel/vg/VecgeomTrackView.hh b/src/geocel/vg/VecgeomTrackView.hh index b9de7c3ae3..1e42ccfc2a 100644 --- a/src/geocel/vg/VecgeomTrackView.hh +++ b/src/geocel/vg/VecgeomTrackView.hh @@ -305,7 +305,7 @@ VecgeomTrackView::operator=(Initializer_t const& init) { #if !CELER_DEVICE_COMPILE auto msg = CELER_LOG_LOCAL(error); - msg << "Failed to initialize geometry state at " << repr(pos_) + msg << "Failed to initialize geometry state at " << repr(pos_) << ' ' << lengthunits::label; #endif failed_ = true; @@ -479,6 +479,23 @@ CELER_FUNCTION Propagation VecgeomTrackView::find_next_step(real_type max_step) *next_surf_ #endif ); + + if (CELER_UNLIKELY(!(next_step_ > 0))) + { +#if !CELER_DEVICE_COMPILE + auto msg = CELER_LOG_LOCAL(error); + msg << "Failed to find next step at " << repr(pos_) << ' ' + << lengthunits::label << " along " << repr(dir_) + << ": computed step is " << repr(next_step_) << ' ' + << lengthunits::label; +#endif + failed_ = true; + Propagation result; + result.distance = next_step_; + result.boundary = false; + return result; + } + if constexpr (CELERITAS_VECGEOM_SURFACE) { // Our accessor uses the next_surf_ state, but the temporary used for diff --git a/test/celeritas/field/FieldPropagator.test.cc b/test/celeritas/field/FieldPropagator.test.cc index 0399af3d7f..af01a6cbcf 100644 --- a/test/celeritas/field/FieldPropagator.test.cc +++ b/test/celeritas/field/FieldPropagator.test.cc @@ -1491,12 +1491,14 @@ TEST_F(CmseTest, coarse) else if (using_solids_vg) { // Bumped (platform-dependent!): counts change a bit - expected_num_boundary[1] = 101; - expected_num_step[1] = 6462; - expected_num_intercept[1] = 19551; - expected_num_integration[1] = 58282; + expected_num_boundary[1] = 37; + expected_num_step[1] = 179; + expected_num_intercept[1] = 614; + expected_num_integration[1] = 1666; static char const* const expected_log_messages[] = { - R"(Moved internally from boundary but safety didn't increase: volume 18 from {10.32, -6.565, 796.9} to {10.32, -6.565, 796.9} (distance: 1.000e-4))"}; + R"(Moved internally from boundary but safety didn't increase: volume 18 from {10.32, -6.565, 796.9} to {10.32, -6.565, 796.9} (distance: 1.000e-4))", + R"(Failed to find next step at {10.32, -6.565, 796.9} cm along {0.6896, -0.1485, 0.7088}: computed step is 0 cm)", + }; EXPECT_VEC_EQ(expected_log_messages, scoped_log_.messages()) << scoped_log_; } diff --git a/test/celeritas/geo/HeuristicGeoExecutor.hh b/test/celeritas/geo/HeuristicGeoExecutor.hh index 39718fdfb0..1ce95554ae 100644 --- a/test/celeritas/geo/HeuristicGeoExecutor.hh +++ b/test/celeritas/geo/HeuristicGeoExecutor.hh @@ -148,7 +148,12 @@ CELER_FUNCTION void HeuristicGeoExecutor::operator()(TrackSlotId tid) const } } - if (prop.boundary) + if (geo.failed()) + { + state.status[tid] = LifeStatus::dead; + return; + } + else if (prop.boundary) { geo.move_to_boundary(); CELER_ASSERT(geo.is_on_boundary()); diff --git a/test/celeritas/global/Stepper.test.cc b/test/celeritas/global/Stepper.test.cc index 0980a02993..7ae369f86a 100644 --- a/test/celeritas/global/Stepper.test.cc +++ b/test/celeritas/global/Stepper.test.cc @@ -189,12 +189,17 @@ TEST_F(SimpleComptonTest, fail_initialize) CELER_TRY_HANDLE(step(make_span(primaries)), LogContextException{this->output_reg().get()}); - static char const* const expected_log_messages[] = { - "Track started outside the geometry", - R"(Killing track {"geo":{"dir":[1.0,0.0,0.0],"is_on_boundary":false,"is_outside":true,"pos":[[1001.,0.0,0.0],"cm"]},"mat":null,"particle":{"energy":[100.0,"MeV"],"particle_id":"gamma"},"sim":{"event_id":0,"num_steps":0,"parent_id":null,"post_step_action":"tracking-cut","status":"errored","step_length":[0.0,"cm"],"time":[0.0,"s"],"track_id":15},"thread_id":31,"track_slot_id":31}: depositing 100 MeV)", - }; if (CELERITAS_UNITS == CELERITAS_UNITS_CGS) { + std::vector expected_log_messages = { + "Track started outside the geometry", + R"(Killing track {"geo":{"dir":[1.0,0.0,0.0],"is_on_boundary":false,"is_outside":true,"pos":[[1001.,0.0,0.0],"cm"]},"mat":null,"particle":{"energy":[100.0,"MeV"],"particle_id":"gamma"},"sim":{"event_id":0,"num_steps":0,"parent_id":null,"post_step_action":"tracking-cut","status":"errored","step_length":[0.0,"cm"],"time":[0.0,"s"],"track_id":15},"thread_id":31,"track_slot_id":31}: depositing 100 MeV)", + }; + if (CELERITAS_CORE_GEO == CELERITAS_CORE_GEO_VECGEOM) + { + expected_log_messages[0] + = "Failed to initialize geometry state at {1001, 0, 0} cm"; + } EXPECT_VEC_EQ(expected_log_messages, scoped_log.messages()); } static char const* const expected_log_levels[] = {"error", "error"}; From 5bcf37793ab6db3cb7c785f073cc4531186efe45 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Fri, 6 Feb 2026 15:06:03 -0500 Subject: [PATCH 12/18] Limit assertion to vecgeom 1 for now --- scripts/cmake-presets/executor.json | 18 +++++------------- src/geocel/vg/VecgeomTrackView.hh | 2 +- 2 files changed, 6 insertions(+), 14 deletions(-) diff --git a/scripts/cmake-presets/executor.json b/scripts/cmake-presets/executor.json index 979172da86..8be814669e 100644 --- a/scripts/cmake-presets/executor.json +++ b/scripts/cmake-presets/executor.json @@ -124,7 +124,7 @@ }, { "name": "vglocal", - "displayName": "With local vecgeom", + "displayName": "With local vecgeom (solids)", "inherits": [".ccache", ".base", ".debug", "default"], "environment": { "CMAKE_PREFIX_PATH": "/opt/spack/opt/spack/tahoe/geant4/11.3.2/gykadvh:/opt/spack/opt/spack/tahoe/covfie/0.15.4/r5g2psb:/opt/spack/opt/spack/tahoe/nlohmann-json/3.12.0/7mlcviu:/opt/spack/opt/spack/tahoe/cli11/2.6.1/znuwv76:/opt/spack/opt/spack/tahoe/googletest/1.17.0/rmxgcak" @@ -137,24 +137,16 @@ "CELERITAS_USE_HepMC3": {"type": "BOOL", "value": "OFF"}, "CELERITAS_USE_ROOT": {"type": "BOOL", "value": "OFF"}, "CELERITAS_USE_VecGeom": {"type": "BOOL", "value": "ON"}, - "CELERITAS_VecGeom_SURFACE": {"type": "BOOL", "value": "ON"}, + "CELERITAS_VecGeom_SURFACE": {"type": "BOOL", "value": "OFF"}, "VecGeom_DIR": {"type": "PATH", "value": "/Users/seth/Code/vecgeom/install/lib/cmake/VecGeom"} } }, { - "name": "vglocal-solid", - "displayName": "With local vecgeom, surfaces disabled", + "name": "vglocal-surface", + "displayName": "With local vecgeom, using surfaces", "inherits": "vglocal", "cacheVariables": { - "CELERITAS_VecGeom_SURFACE": {"type": "BOOL", "value": "OFF"} - } - }, - { - "name": "vglocal-v1", - "displayName": "With local vecgeom, surfaces disabled", - "inherits": "vglocal", - "cacheVariables": { - "VecGeom_DIR": {"type": "PATH", "value": "/Users/seth/Code/vecgeom-v1/install/lib/cmake/VecGeom"} + "CELERITAS_VecGeom_SURFACE": {"type": "BOOL", "value": "ON"} } }, { diff --git a/src/geocel/vg/VecgeomTrackView.hh b/src/geocel/vg/VecgeomTrackView.hh index 1e42ccfc2a..c6e8defff4 100644 --- a/src/geocel/vg/VecgeomTrackView.hh +++ b/src/geocel/vg/VecgeomTrackView.hh @@ -480,7 +480,7 @@ CELER_FUNCTION Propagation VecgeomTrackView::find_next_step(real_type max_step) #endif ); - if (CELER_UNLIKELY(!(next_step_ > 0))) + if (CELER_UNLIKELY(VECGEOM_VERSION < 0x020000 && !(next_step_ > 0))) { #if !CELER_DEVICE_COMPILE auto msg = CELER_LOG_LOCAL(error); From d9bdca847808f15dbe54a3d0131c0ebc5d703b93 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Fri, 6 Feb 2026 19:57:35 -0500 Subject: [PATCH 13/18] Restore field propagator result for vg 2 solid and delete unchecked g4 result --- test/celeritas/field/FieldPropagator.test.cc | 23 ++++++++++++-------- 1 file changed, 14 insertions(+), 9 deletions(-) diff --git a/test/celeritas/field/FieldPropagator.test.cc b/test/celeritas/field/FieldPropagator.test.cc index af01a6cbcf..916afe578b 100644 --- a/test/celeritas/field/FieldPropagator.test.cc +++ b/test/celeritas/field/FieldPropagator.test.cc @@ -1472,15 +1472,7 @@ TEST_F(CmseTest, coarse) std::vector expected_num_integration = {80659, 58204, 41914, 26114}; std::vector expected_log_messages; - if (CELERITAS_CORE_GEO == CELERITAS_CORE_GEO_GEANT4) - { - // FIXME: this happens because of incorrect momentum update - expected_num_boundary = {134, 37, 60, 40}; - expected_num_step = {10001, 179, 3236, 1303}; - expected_num_intercept = {30419, 615, 16170, 9956}; - expected_num_integration = {80659, 1670, 41914, 26114}; - } - else if (using_surface_vg) + if (using_surface_vg) { expected_num_boundary = {134, 37, 43, 16}; expected_num_step = {10001, 179, 160, 63}; @@ -1488,6 +1480,19 @@ TEST_F(CmseTest, coarse) expected_num_integration = {80659, 1670, 1956, 1092}; EXPECT_TRUE(scoped_log_.empty()) << scoped_log_; } + else if (using_solids_vg && CELERITAS_VECGEOM_VERSION >= 0x020000) + { + // Bumped (platform-dependent!): counts change a bit + expected_num_boundary[1] = 101; + expected_num_step[1] = 6462; + expected_num_intercept[1] = 19551; + expected_num_integration[1] = 58282; + static char const* const expected_log_messages[] = { + R"(Moved internally from boundary but safety didn't increase: volume 18 from {10.32, -6.565, 796.9} to {10.32, -6.565, 796.9} (distance: 1.000e-4))", + }; + EXPECT_VEC_EQ(expected_log_messages, scoped_log_.messages()) + << scoped_log_; + } else if (using_solids_vg) { // Bumped (platform-dependent!): counts change a bit From db83ec1aa3bf5856079920184956dcd3a7b9225c Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Fri, 6 Feb 2026 20:13:41 -0500 Subject: [PATCH 14/18] Fail if boundary is not crossed --- src/geocel/vg/VecgeomTrackView.hh | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/src/geocel/vg/VecgeomTrackView.hh b/src/geocel/vg/VecgeomTrackView.hh index c6e8defff4..6c1270cc9e 100644 --- a/src/geocel/vg/VecgeomTrackView.hh +++ b/src/geocel/vg/VecgeomTrackView.hh @@ -597,6 +597,19 @@ CELER_FUNCTION void VecgeomTrackView::cross_boundary() vgnext_); } + // Relocation should have changed volume + if (CELER_UNLIKELY(vgstate_.HasSamePathAsOther(vgnext_))) + { +#if !CELER_DEVICE_COMPILE + auto msg = CELER_LOG_LOCAL(error); + msg << "Failed to cross boundary: unique volume instance is the same " + "before and after at " + << repr(pos_) << ' ' << lengthunits::label; +#endif + failed_ = true; + return; + } + vgstate_ = vgnext_; CELER_ENSURE(this->is_on_boundary()); From b84a2449264d3599209adae2b10739ee572d03a3 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Fri, 6 Feb 2026 20:13:44 -0500 Subject: [PATCH 15/18] Revert "Fail if boundary is not crossed" This reverts commit db83ec1aa3bf5856079920184956dcd3a7b9225c. --- src/geocel/vg/VecgeomTrackView.hh | 13 ------------- 1 file changed, 13 deletions(-) diff --git a/src/geocel/vg/VecgeomTrackView.hh b/src/geocel/vg/VecgeomTrackView.hh index 6c1270cc9e..c6e8defff4 100644 --- a/src/geocel/vg/VecgeomTrackView.hh +++ b/src/geocel/vg/VecgeomTrackView.hh @@ -597,19 +597,6 @@ CELER_FUNCTION void VecgeomTrackView::cross_boundary() vgnext_); } - // Relocation should have changed volume - if (CELER_UNLIKELY(vgstate_.HasSamePathAsOther(vgnext_))) - { -#if !CELER_DEVICE_COMPILE - auto msg = CELER_LOG_LOCAL(error); - msg << "Failed to cross boundary: unique volume instance is the same " - "before and after at " - << repr(pos_) << ' ' << lengthunits::label; -#endif - failed_ = true; - return; - } - vgstate_ = vgnext_; CELER_ENSURE(this->is_on_boundary()); From be6fb626f8a1d584996c3b59e6b41bc0eb6f699d Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Sat, 7 Feb 2026 07:41:30 -0500 Subject: [PATCH 16/18] Loosen distance requirements to allow geometry to limit step internally --- src/celeritas/field/FieldPropagator.hh | 3 ++- src/celeritas/field/LinearPropagator.hh | 5 ++++- test/celeritas/geo/HeuristicGeoExecutor.hh | 2 +- 3 files changed, 7 insertions(+), 3 deletions(-) diff --git a/src/celeritas/field/FieldPropagator.hh b/src/celeritas/field/FieldPropagator.hh index 49d9f4f272..f8c3ff601a 100644 --- a/src/celeritas/field/FieldPropagator.hh +++ b/src/celeritas/field/FieldPropagator.hh @@ -294,6 +294,7 @@ FieldPropagator::operator()(real_type step) -> result_type // what step length we took, which means we're stuck. // Using the just-reapplied direction, hope that we're pointing deeper // into the current volume and bump the particle. + // TODO: move this into a higher-level failure/bump mechanic result.distance = celeritas::min(this->bump_distance(), step); result.boundary = false; axpy(result.distance, dir, &state_.pos); @@ -308,7 +309,7 @@ FieldPropagator::operator()(real_type step) -> result_type // within the driver, the distance may be very slightly beyond the // requested step. CELER_ENSURE( - result.distance > 0 + (result.distance > 0 || geo_.failed()) && (result.distance <= step || soft_equal(result.distance, step))); return result; } diff --git a/src/celeritas/field/LinearPropagator.hh b/src/celeritas/field/LinearPropagator.hh index 7463144570..8fd39ea5ca 100644 --- a/src/celeritas/field/LinearPropagator.hh +++ b/src/celeritas/field/LinearPropagator.hh @@ -62,7 +62,10 @@ CELER_FUNCTION auto LinearPropagator::operator()(real_type dist) } else { - CELER_ASSERT(dist == result.distance); + // TODO: allow distance traveled to be less than the input distance for + // cases like bumps, internal boundary, etc. + CELER_ASSERT(result.distance <= dist + && (result.distance > 0 || geo_.failed())); geo_.move_internal(dist); } diff --git a/test/celeritas/geo/HeuristicGeoExecutor.hh b/test/celeritas/geo/HeuristicGeoExecutor.hh index 1ce95554ae..6eb81ac87a 100644 --- a/test/celeritas/geo/HeuristicGeoExecutor.hh +++ b/test/celeritas/geo/HeuristicGeoExecutor.hh @@ -162,7 +162,7 @@ CELER_FUNCTION void HeuristicGeoExecutor::operator()(TrackSlotId tid) const { // Check for similar assertions in FieldPropagator before loosening // this one! - CELER_ASSERT(prop.distance == step); + CELER_ASSERT(prop.distance <= step); CELER_ASSERT(prop.distance > 0); #if CELERITAS_DEBUG auto orig_pos = geo.pos(); From 53e4f38f273ebba0540c55a21569652357b2e8a6 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Sat, 7 Feb 2026 07:41:39 -0500 Subject: [PATCH 17/18] Clear result in case of nan --- src/geocel/vg/VecgeomTrackView.hh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/geocel/vg/VecgeomTrackView.hh b/src/geocel/vg/VecgeomTrackView.hh index c6e8defff4..4569d2943b 100644 --- a/src/geocel/vg/VecgeomTrackView.hh +++ b/src/geocel/vg/VecgeomTrackView.hh @@ -491,7 +491,7 @@ CELER_FUNCTION Propagation VecgeomTrackView::find_next_step(real_type max_step) #endif failed_ = true; Propagation result; - result.distance = next_step_; + result.distance = 0; result.boundary = false; return result; } From 6ce0b814aebd632c4fea2fb67708caaac6856610 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Sat, 7 Feb 2026 07:41:50 -0500 Subject: [PATCH 18/18] Try to work around spurious CUDA error --- src/geocel/vg/detail/VecgeomSetup.cu | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/src/geocel/vg/detail/VecgeomSetup.cu b/src/geocel/vg/detail/VecgeomSetup.cu index 471cb1e713..2fb4795d18 100644 --- a/src/geocel/vg/detail/VecgeomSetup.cu +++ b/src/geocel/vg/detail/VecgeomSetup.cu @@ -191,14 +191,14 @@ CudaPointers navindex_pointers_device() void check_other_device_pointers() { - if constexpr (VECGEOM_VERSION < 0x020000) - return; - - CELER_VALIDATE(get_device_pointer() != nullptr, - << "failed to copy VG logical volumes to GPU"); + if constexpr (VECGEOM_VERSION >= 0x020000) + { + CELER_VALIDATE(get_device_pointer() != nullptr, + << "failed to copy VG logical volumes to GPU"); - CELER_VALIDATE(get_device_pointer() != nullptr, - << "failed to copy VG places volumes to GPU"); + CELER_VALIDATE(get_device_pointer() != nullptr, + << "failed to copy VG places volumes to GPU"); + } } #if CELER_VGNAV == CELER_VGNAV_TUPLE