VecGeom/navigation/NavStateTuple.h

0001 /// \file NavStateTuple.h
0002 /// \author Andrei Gheata (andrei.gheata@cern.ch)
0003 /// \date 20.06.2023
0004
0005 #ifndef VECGEOM_NAVIGATION_NAVSTATETUPLE_H_
0006 #define VECGEOM_NAVIGATION_NAVSTATETUPLE_H_
0007
0008 #include "VecGeom/base/Config.h"
0009 #include "VecGeom/base/Global.h"
0010 #include "VecGeom/base/Transformation3D.h"
0011 #include "VecGeom/base/Transformation3DMP.h"
0012 #include "VecGeom/volumes/PlacedVolume.h"
0013 #include "VecGeom/volumes/VolumeTree.h"
0014 #include "VecGeom/management/GeoManager.h"
0015 #include "VecGeom/management/DeviceGlobals.h"
0016
0017 #include <iostream>
0018 #include <list>
0019 #include <sstream>
0020
0021 namespace vecgeom {
0022
0023 template <uint MAX_DEPTH>
0024 struct NavTuple {
0025   NavIndex_t fNavInd[MAX_DEPTH]{0};
0026   uint fLevel{0};
0027
0028   NavTuple() = default;
0029
0030   VECCORE_ATT_HOST_DEVICE
0031   NavTuple(NavTuple<MAX_DEPTH> const &other) : fLevel{other.fLevel}
0032   {
0033     for (uint i = 0; i <= fLevel; ++i)
0034       fNavInd[i] = other.fNavInd[i];
0035   }
0036
0037   VECCORE_ATT_HOST_DEVICE
0038   NavTuple(NavIndex_t ind) { fNavInd[0] = ind; }
0039
0040   template <typename Container>
0041   VECCORE_ATT_HOST_DEVICE NavTuple(Container const *cont)
0042   {
0043     VECGEOM_ASSERT(cont->size() <= MAX_DEPTH);
0044     for (NavIndex_t navind : *cont)
0045       fNavInd[fLevel++] = navind;
0046     if (fLevel > 0) fLevel--;
0047   }
0048
0049   VECGEOM_FORCE_INLINE
0050   VECCORE_ATT_HOST_DEVICE
0051   static constexpr uint GetMaxDepth() { return MAX_DEPTH; }
0052
0053   VECGEOM_FORCE_INLINE
0054   VECCORE_ATT_HOST_DEVICE
0055   NavTuple<MAX_DEPTH> &operator=(NavTuple<MAX_DEPTH> const &other)
0056   {
0057     fLevel = other.fLevel;
0058     for (uint i = 0; i <= fLevel; ++i)
0059       fNavInd[i] = other.fNavInd[i];
0060     return *this;
0061   }
0062
0063   /// @brief Assign a single-level navigation index
0064   /// @param ind Navigation index at first level
0065   VECGEOM_FORCE_INLINE
0066   VECCORE_ATT_HOST_DEVICE
0067   NavTuple<MAX_DEPTH> &operator=(NavIndex_t ind)
0068   {
0069     fLevel     = 0;
0070     fNavInd[0] = ind;
0071     return *this;
0072   }
0073
0074   VECGEOM_FORCE_INLINE
0075   VECCORE_ATT_HOST_DEVICE
0076   NavIndex_t operator[](uint i) const
0077   {
0078     VECGEOM_VALIDATE(i < MAX_DEPTH, << "NavTuple::operator[] out of range");
0079     return (i < MAX_DEPTH) ? fNavInd[i] : 0;
0080   }
0081
0082   VECGEOM_FORCE_INLINE
0083   VECCORE_ATT_HOST_DEVICE
0084   NavIndex_t &operator[](uint i) { return fNavInd[i]; }
0085
0086   VECGEOM_FORCE_INLINE
0087   VECCORE_ATT_HOST_DEVICE
0088   bool operator==(NavTuple<MAX_DEPTH> const &other) const
0089   {
0090     if (fLevel != other.fLevel) return false;
0091     for (int i = fLevel; i >= 0; i--) {
0092       if (fNavInd[i] != other.fNavInd[i]) return false;
0093     }
0094     return true;
0095   }
0096
0097   VECGEOM_FORCE_INLINE
0098   VECCORE_ATT_HOST_DEVICE
0099   bool operator!=(NavTuple<MAX_DEPTH> const &other) const { return !operator==(other); }
0100
0101   VECGEOM_FORCE_INLINE
0102   VECCORE_ATT_HOST_DEVICE
0103   bool operator<(NavTuple<MAX_DEPTH> const &other) const
0104   {
0105     if (fLevel == other.fLevel) {
0106       for (unsigned i = 0; i <= fLevel; ++i) {
0107         if (fNavInd[i] < other.fNavInd[i]) return true;
0108         if (fNavInd[i] > other.fNavInd[i]) return false;
0109       }
0110     } else {
0111       return fLevel < other.fLevel;
0112     }
0113     return false;
0114   }
0115
0116   VECGEOM_FORCE_INLINE
0117   VECCORE_ATT_HOST_DEVICE
0118   bool operator==(NavIndex_t navind) const { return fLevel == 0 && fNavInd[0] == navind; }
0119
0120   VECGEOM_FORCE_INLINE
0121   VECCORE_ATT_HOST_DEVICE
0122   bool operator!=(NavIndex_t navind) const { return !operator==(navind); }
0123
0124   VECGEOM_FORCE_INLINE
0125   VECCORE_ATT_HOST_DEVICE
0126   void Clear()
0127   {
0128     fLevel     = 0;
0129     fNavInd[0] = 0;
0130   }
0131
0132   VECGEOM_FORCE_INLINE
0133   VECCORE_ATT_HOST_DEVICE
0134   bool IsOutside() const { return (fLevel == 0) && (fNavInd[0] == 0); }
0135
0136   VECGEOM_FORCE_INLINE
0137   VECCORE_ATT_HOST_DEVICE
0138   void Push(NavIndex_t value)
0139   {
0140     if (!IsOutside()) fLevel++;
0141     VECGEOM_VALIDATE(fLevel < MAX_DEPTH, << "NavTuple::Push out of range");
0142     if (fLevel < MAX_DEPTH) fNavInd[fLevel] = value;
0143   }
0144
0145   VECGEOM_FORCE_INLINE
0146   VECCORE_ATT_HOST_DEVICE
0147   void Set(NavIndex_t value)
0148   {
0149     VECGEOM_VALIDATE(fLevel < MAX_DEPTH, << "NavTuple::Set out of range");
0150     if (fLevel < MAX_DEPTH) fNavInd[fLevel] = value;
0151   }
0152
0153   VECGEOM_FORCE_INLINE
0154   VECCORE_ATT_HOST_DEVICE
0155   NavIndex_t Top() const
0156   {
0157     VECGEOM_VALIDATE(fLevel < MAX_DEPTH, << "NavTuple::Top out of range");
0158     return (fLevel < MAX_DEPTH) ? fNavInd[fLevel] : 0;
0159   }
0160 };
0161
0162 template <unsigned int MAX_DEPTH>
0163 std::ostream &operator<<(std::ostream &os, NavTuple<MAX_DEPTH> const &nav_tuple)
0164 {
0165   os << "(" << nav_tuple[0];
0166   for (unsigned i = 1; i < nav_tuple.fLevel; ++i)
0167     os << " ," << nav_tuple[i];
0168   os << ")";
0169   return os;
0170 }
0171
0172 using NavTuple_t = NavTuple<VECGEOM_NAVTUPLE_MAXDEPTH>;
0173
0174 /**
0175  * @brief A class describing a current geometry state based on a tuple of indices
0176  */
0177 class NavStateTuple {
0178 public:
0179   using Value_t = NavTuple_t;
0180
0181 private:
0182   NavTuple_t fNavTuple{0};   ///< Navigation state tuple
0183   NavTuple_t fLastExited{0}; ///< Navigation state index of the last exited state
0184   bool fOnBoundary = false;  ///< flag indicating whether track is on boundary of the "Top()" placed volume
0185
0186 public:
0187   VECCORE_ATT_HOST_DEVICE
0188   NavStateTuple(NavTuple_t nav_tpl = 0) : fNavTuple(nav_tpl) {}
0189
0190   template <typename Container>
0191   VECCORE_ATT_HOST_DEVICE NavStateTuple(Container const *cont) : fNavTuple(cont)
0192   {
0193   }
0194
0195   VECGEOM_FORCE_INLINE
0196   VECCORE_ATT_HOST_DEVICE
0197   static unsigned char GetMaxLevel()
0198   {
0199 #ifdef VECCORE_CUDA_DEVICE_COMPILATION
0200     return vecgeom::globaldevicegeomdata::gMaxDepth;
0201 #else
0202     return (unsigned char)GeoManager::Instance().getMaxDepth();
0203 #endif
0204   }
0205
0206   // Static accessors
0207   VECCORE_ATT_HOST_DEVICE
0208   static NavStateTuple *MakeInstance(int)
0209   {
0210     // MaxLevel is 'zero' based (i.e. maxlevel==0 requires one value)
0211     return new NavStateTuple();
0212   }
0213
0214   VECCORE_ATT_HOST_DEVICE
0215   static NavStateTuple *MakeCopy(NavStateTuple const &other) { return new NavStateTuple(other); }
0216
0217   VECCORE_ATT_HOST_DEVICE
0218   static NavStateTuple *MakeInstanceAt(int, void *addr) { return new (addr) NavStateTuple(); }
0219
0220   VECCORE_ATT_HOST_DEVICE
0221   static NavStateTuple *MakeCopy(NavStateTuple const &other, void *addr) { return new (addr) NavStateTuple(other); }
0222
0223   VECCORE_ATT_HOST_DEVICE
0224   static void ReleaseInstance(NavStateTuple *state)
0225   {
0226     // MaxLevel is 'zero' based (i.e. maxlevel==0 requires one value)
0227     delete state;
0228   }
0229
0230   // returns the size in bytes of a NavStateTuple object with internal
0231   // path depth maxlevel
0232   VECCORE_ATT_HOST_DEVICE
0233   static size_t SizeOfInstance(int)
0234   {
0235     // MaxLevel is 'zero' based (i.e. maxlevel==0 requires one value)
0236     return sizeof(NavStateTuple);
0237   }
0238
0239   // returns the size in bytes of a NavStateIndex object with internal
0240   // path depth maxlevel -- including space needed for padding to next aligned object
0241   // of same kind
0242   VECCORE_ATT_HOST_DEVICE
0243   static size_t SizeOfInstanceAlignAware(int)
0244   {
0245     // MaxLevel is 'zero' based (i.e. maxlevel==0 requires one value)
0246     return sizeof(NavStateTuple);
0247   }
0248
0249   VECCORE_ATT_HOST_DEVICE
0250   VECGEOM_FORCE_INLINE
0251   NavIndex_t GetNavIndex() const { return fNavTuple.Top(); }
0252
0253   VECCORE_ATT_HOST_DEVICE
0254   VECGEOM_FORCE_INLINE
0255   NavTuple_t const &GetState() const { return fNavTuple; }
0256
0257   VECCORE_ATT_HOST_DEVICE
0258   int GetObjectSize() const { return (int)sizeof(NavStateTuple); }
0259
0260   VECCORE_ATT_HOST_DEVICE
0261   static size_t SizeOf(size_t) { return sizeof(NavStateTuple); }
0262
0263   VECCORE_ATT_HOST_DEVICE
0264   void CopyTo(NavStateTuple *other) const { *other = *this; }
0265
0266   // copies a fixed and predetermined number of bytes
0267   // might be useful for specialized navigators which know the depth + SizeOf in advance
0268   // N is number of bytes to be copied and can be obtained by a prior call to constexpr NavStateIndex::SizeOf( ... );
0269   template <size_t N>
0270   void CopyToFixedSize(NavStateTuple *other) const
0271   {
0272     *other = *this;
0273   }
0274
0275   /// @brief Returns the navigation index address in the table. This is needed for calculating
0276   //         the addresses of components associated with nav_ind
0277   /// @param nav_ind Navigation index
0278   /// @return Start address fir the meta information associated with the index
0279   VECCORE_ATT_HOST_DEVICE
0280   VECGEOM_FORCE_INLINE
0281   static NavIndex_t const *NavIndAddr(NavIndex_t nav_ind)
0282   {
0283 #ifdef VECCORE_CUDA_DEVICE_COMPILATION
0284     // checking here for NVCC_DEVICE since the global variable globaldevicegeomgata::gCompact...
0285     // is marked __device__ and can only be compiled within device compiler passes
0286     VECGEOM_ASSERT(vecgeom::globaldevicegeomdata::gNavIndex != nullptr);
0287     return &vecgeom::globaldevicegeomdata::gNavIndex[nav_ind];
0288 #else
0289     VECGEOM_ASSERT(vecgeom::GeoManager::gNavIndex != nullptr);
0290     return &vecgeom::GeoManager::gNavIndex[nav_ind];
0291 #endif
0292   }
0293
0294   /// @brief Returns the navigation index stored at a given index in the table
0295   /// @param i Index in the navigation table
0296   /// @return Navigation index stored at the requested index
0297   VECCORE_ATT_HOST_DEVICE
0298   VECGEOM_FORCE_INLINE
0299   static NavIndex_t NavInd(NavIndex_t i) { return *NavIndAddr(i); }
0300
0301   VECCORE_ATT_HOST_DEVICE
0302   VECGEOM_FORCE_INLINE
0303   static VPlacedVolume const *ToPlacedVolume(size_t index)
0304   {
0305 #ifdef VECCORE_CUDA_DEVICE_COMPILATION
0306     // checking here for NVCC_DEVICE since the global variable globaldevicegeomgata::gCompact...
0307     // is marked __device__ and can only be compiled within device compiler passes
0308     VECGEOM_ASSERT(vecgeom::globaldevicegeomdata::gCompactPlacedVolBuffer != nullptr);
0309     return &vecgeom::globaldevicegeomdata::gCompactPlacedVolBuffer[index];
0310 #else
0311     VECGEOM_ASSERT(vecgeom::GeoManager::gCompactPlacedVolBuffer == nullptr ||
0312                    vecgeom::GeoManager::gCompactPlacedVolBuffer[index].id() == index);
0313     return &vecgeom::GeoManager::gCompactPlacedVolBuffer[index];
0314 #endif
0315   }
0316
0317   VECCORE_ATT_HOST_DEVICE
0318   VECGEOM_FORCE_INLINE
0319   static int WorldId() { return NavInd(2); }
0320
0321   VECCORE_ATT_HOST_DEVICE
0322   VECGEOM_FORCE_INLINE
0323   static vecgeom::PlacedId const &ToPlacedId(size_t iplaced)
0324   {
0325     return vecgeom::VolumeTree::Instance().fPlaced[iplaced];
0326   }
0327
0328   /// @brief Implementation for getting the logical id for a given navigation index
0329   VECCORE_ATT_HOST_DEVICE
0330   VECGEOM_FORCE_INLINE
0331   static unsigned int GetLogicalIdImpl(NavIndex_t nav_index) { return nav_index ? NavInd(NavInd(nav_index + 4)) : 0; }
0332
0333   /// @brief Implementation for getting the logical id for a given navigation tuple
0334   VECCORE_ATT_HOST_DEVICE
0335   VECGEOM_FORCE_INLINE
0336   static unsigned int GetLogicalIdImpl(NavTuple_t const &nav_tuple) { return GetLogicalIdImpl(nav_tuple.Top()); }
0337
0338   /// @brief Implementation for getting the child id for a given navigation index
0339   VECCORE_ATT_HOST_DEVICE
0340   VECGEOM_FORCE_INLINE
0341   static int GetChildIdImpl(NavIndex_t const &nav_index)
0342   {
0343     auto content_ichild = reinterpret_cast<const int *>(NavIndAddr(nav_index + 2));
0344     return *content_ichild;
0345   }
0346
0347   /// @brief Implementation for getting the child id for a given navigation tuple
0348   VECCORE_ATT_HOST_DEVICE
0349   VECGEOM_FORCE_INLINE
0350   static int GetChildIdImpl(NavTuple_t const &nav_tuple) { return GetChildIdImpl(nav_tuple.Top()); }
0351
0352   /// @brief Implementation for getting the number of daughters for a given navigation index
0353   VECCORE_ATT_HOST_DEVICE
0354   VECGEOM_FORCE_INLINE
0355   static unsigned int GetNdaughtersImpl(NavIndex_t const &nav_index) { return NavInd(NavInd(nav_index + 4) + 1); }
0356
0357   /// @brief Implementation for getting the number of daughters for a given navigation tuple
0358   VECCORE_ATT_HOST_DEVICE
0359   VECGEOM_FORCE_INLINE
0360   static unsigned int GetNdaughtersImpl(NavTuple_t const &nav_tuple) { return NavInd(NavInd(nav_tuple.Top() + 4) + 1); }
0361
0362   /// @brief Implementation for getting the scene id for a given navigation index
0363   VECCORE_ATT_HOST_DEVICE
0364   VECGEOM_FORCE_INLINE
0365   static bool GetSceneIdImpl(NavIndex_t nav_ind, unsigned short &scene_id, unsigned short &newscene_id)
0366   {
0367     scene_id    = 0;
0368     newscene_id = 0;
0369     if (nav_ind == 0) return false;
0370     auto scenes = reinterpret_cast<const unsigned short *>(NavIndAddr(nav_ind + 5));
0371     scene_id    = scenes[0];
0372     newscene_id = scenes[1];
0373     return (newscene_id != scene_id);
0374   }
0375
0376   /// @brief Implementation for getting the scene id for a given navigation tuple
0377   /// @return True if the state points to a scene volume
0378   VECCORE_ATT_HOST_DEVICE
0379   VECGEOM_FORCE_INLINE
0380   static bool GetSceneIdImpl(NavTuple_t const &nav_tuple, unsigned short &scene_id, unsigned short &newscene_id)
0381   {
0382     auto top = nav_tuple.Top();
0383     if (top == 0) {
0384       scene_id    = GetParentNewSceneImpl(nav_tuple);
0385       newscene_id = scene_id;
0386       return true;
0387     }
0388     return GetSceneIdImpl(nav_tuple.Top(), scene_id, newscene_id);
0389   }
0390
0391   /// @brief Implementation for getting if a given tuple points to a scene volume
0392   VECCORE_ATT_HOST_DEVICE
0393   VECGEOM_FORCE_INLINE
0394   static bool IsSceneImpl(NavTuple_t const &nav_tuple)
0395   {
0396     unsigned short scene_id = 0, newscene_id = 0;
0397     return GetSceneIdImpl(nav_tuple.Top(), scene_id, newscene_id);
0398   }
0399
0400   /// @brief Implementation for getting the scene level for a given navigation tuple
0401   VECCORE_ATT_HOST_DEVICE
0402   VECGEOM_FORCE_INLINE
0403   static unsigned int GetSceneLevelImpl(NavTuple_t const &nav_tuple) { return nav_tuple.fLevel; }
0404
0405   /// @brief Implementation for getting the parent scene id
0406   VECCORE_ATT_HOST_DEVICE
0407   VECGEOM_FORCE_INLINE
0408   static unsigned short GetParentSceneImpl(NavTuple_t const &nav_tuple)
0409   {
0410     if (nav_tuple.fLevel < 1) return 0;
0411     auto top_parent = nav_tuple[nav_tuple.fLevel - 1];
0412     return *reinterpret_cast<const unsigned short *>(NavIndAddr(top_parent + 5));
0413   }
0414
0415   /// @brief Implementation for getting the parent scene id
0416   VECCORE_ATT_HOST_DEVICE
0417   VECGEOM_FORCE_INLINE
0418   static unsigned short GetParentNewSceneImpl(NavTuple_t const &nav_tuple)
0419   {
0420     if (nav_tuple.fLevel < 1) return 0;
0421     auto top_parent = nav_tuple[nav_tuple.fLevel - 1];
0422     auto scenes     = reinterpret_cast<const unsigned short *>(NavIndAddr(top_parent + 4));
0423     return scenes[1];
0424   }
0425
0426   /// @brief Implementation for getting the local level associated with a navigation index in the current scene
0427   VECGEOM_FORCE_INLINE
0428   VECCORE_ATT_HOST_DEVICE
0429   static unsigned char GetLevelImpl(NavIndex_t nav_ind)
0430   {
0431     auto content_level = reinterpret_cast<const unsigned char *>(NavIndAddr(nav_ind + 6));
0432     return *content_level;
0433   }
0434
0435   /// @brief Implementation for getting the global level associated with a navigation tuple
0436   VECGEOM_FORCE_INLINE
0437   VECCORE_ATT_HOST_DEVICE
0438   static unsigned char GetLevelImpl(NavTuple_t const &nav_tuple)
0439   {
0440     int level = 0;
0441     for (uint ituple = 0; ituple <= nav_tuple.fLevel; ++ituple) {
0442       level += GetLevelImpl(nav_tuple[ituple]);
0443     }
0444     return level;
0445   }
0446
0447   VECCORE_ATT_HOST_DEVICE
0448   VECGEOM_FORCE_INLINE
0449   static NavTuple_t GetNavTupleImpl(NavTuple_t nav_tuple, int level)
0450   {
0451     int level_max = GetLevelImpl(nav_tuple);
0452     int up        = level_max - level;
0453     VECGEOM_VALIDATE(up >= 0, << "called GetNavIndexImpl for level larger than current level");
0454     NavIndex_t mother = nav_tuple.Top();
0455     while (up--) {
0456       mother = NavInd(mother);
0457       if (mother == 0) {
0458         if (nav_tuple.fLevel == 0) break;
0459         mother = nav_tuple[--nav_tuple.fLevel];
0460       }
0461     }
0462     nav_tuple.Set(mother);
0463     return nav_tuple;
0464   }
0465
0466   VECGEOM_FORCE_INLINE
0467   VECCORE_ATT_HOST_DEVICE
0468   static NavIndex_t GetIdImpl(NavIndex_t nav_ind) { return (nav_ind > 0) ? NavInd(nav_ind + 3) : 0; }
0469
0470   VECGEOM_FORCE_INLINE
0471   VECCORE_ATT_HOST_DEVICE
0472   static bool IsDescendentImpl(NavTuple_t const &child_ind, NavTuple_t const &parent_ind)
0473   {
0474     NavTuple_t ind = child_ind;
0475     while (parent_ind < ind) {
0476       PopImpl(ind);
0477       if (ind == parent_ind) return true;
0478     }
0479     return false;
0480   }
0481
0482   VECGEOM_FORCE_INLINE
0483   VECCORE_ATT_HOST_DEVICE
0484   static bool IsDescendentImpl(NavIndex_t child_ind, NavIndex_t parent_ind)
0485   {
0486     // Only search inside a scene
0487     auto ind = child_ind;
0488     while (ind > parent_ind) {
0489       ind = NavInd(ind);
0490       if (ind == parent_ind) return true;
0491     }
0492     return false;
0493   }
0494
0495   VECGEOM_FORCE_INLINE
0496   VECCORE_ATT_HOST_DEVICE
0497   static NavIndex_t GetChildNavInd(NavIndex_t nav_ind, int ichild)
0498   {
0499     return (nav_ind > 0) ? NavInd(NavInd(nav_ind + 4) + 2 + ichild) : 0;
0500   }
0501
0502   VECCORE_ATT_HOST_DEVICE
0503   static bool IsValid(NavTuple_t const &nav_tuple, int nprint = 0)
0504   {
0505     bool valid = true;
0506     if (nav_tuple.fLevel == 0 && nav_tuple.Top() <= 1) return valid;
0507     for (unsigned i = 0; i <= nav_tuple.fLevel; ++i) {
0508       auto nav_ind = nav_tuple[i];
0509       if (nav_ind == 0) return false;
0510       if (nav_ind == 1) continue;
0511       auto parent = NavInd(nav_ind);
0512       if (parent > 0) {
0513         // Check the pointer to nav_ind in parent record
0514         auto ichild        = GetChildIdImpl(nav_ind);
0515         auto nav_ind_child = GetChildNavInd(parent, ichild);
0516         valid &= nav_ind_child == nav_ind;
0517       }
0518     }
0519     return valid;
0520   }
0521
0522   VECGEOM_FORCE_INLINE
0523   VECCORE_ATT_HOST_DEVICE
0524   static void PopImpl(NavTuple_t &nav_tuple)
0525   {
0526     auto top = nav_tuple.Top();
0527     if (!top) return;
0528     top = NavInd(top);
0529     nav_tuple.Set(top);
0530     if (!top && nav_tuple.fLevel) nav_tuple.fLevel--;
0531   }
0532
0533   VECGEOM_FORCE_INLINE
0534   VECCORE_ATT_HOST_DEVICE
0535   static void PopImpl(NavIndex_t &nav_index)
0536   {
0537     if (!nav_index) return;
0538     nav_index = NavInd(nav_index);
0539   }
0540
0541   VECGEOM_FORCE_INLINE
0542   VECCORE_ATT_HOST_DEVICE
0543   static void PushImpl(NavTuple_t &nav_tuple, VPlacedVolume const *v)
0544   {
0545     auto top = nav_tuple.Top();
0546     if (top) {
0547       auto child       = NavInd(NavInd(top + 4) + 2 + v->GetChildId());
0548       bool on_newscene = NavInd(child) == 0;
0549       if (on_newscene) {
0550         nav_tuple.fLevel++;
0551         VECGEOM_ASSERT(nav_tuple.fLevel < NavTuple_t::GetMaxDepth());
0552       }
0553       nav_tuple.Set(child);
0554     } else {
0555       nav_tuple.Set(1);
0556     }
0557   }
0558
0559   VECGEOM_FORCE_INLINE
0560   VECCORE_ATT_HOST_DEVICE
0561   static void PushDaughterImpl(NavTuple_t &nav_tuple, int idaughter)
0562   {
0563     auto top = nav_tuple.Top();
0564     if (top) {
0565       VECGEOM_ASSERT(idaughter >= 0 && idaughter < int(GetNdaughtersImpl(top)));
0566       auto child       = NavInd(NavInd(top + 4) + 2 + idaughter);
0567       bool on_newscene = NavInd(child) == 0;
0568       if (on_newscene) {
0569         nav_tuple.fLevel++;
0570         VECGEOM_ASSERT(nav_tuple.fLevel < NavTuple_t::GetMaxDepth());
0571       }
0572       nav_tuple.Set(child);
0573     } else {
0574       nav_tuple.Set(1);
0575     }
0576   }
0577
0578   VECGEOM_FORCE_INLINE
0579   VECCORE_ATT_HOST_DEVICE
0580   static void PushImpl(NavTuple_t &nav_tuple, int iplaced)
0581   {
0582     auto const &pv_ind = ToPlacedId(iplaced);
0583     PushDaughterImpl(nav_tuple, pv_ind.fChildId);
0584   }
0585
0586   VECGEOM_FORCE_INLINE
0587   VECCORE_ATT_HOST_DEVICE
0588   static VPlacedVolume const *TopImpl(NavTuple_t const &nav_tuple)
0589   {
0590     auto top = nav_tuple.Top();
0591     return (top > 0) ? ToPlacedVolume(NavInd(top + 1)) : nullptr;
0592   }
0593
0594   VECGEOM_FORCE_INLINE
0595   VECCORE_ATT_HOST_DEVICE
0596   static VPlacedVolume const *World() { return ToPlacedVolume(NavInd(2)); }
0597
0598   VECGEOM_FORCE_INLINE
0599   VECCORE_ATT_HOST_DEVICE
0600   static int TopIdImpl(NavTuple_t const &nav_tuple)
0601   {
0602     auto top = nav_tuple.Top();
0603     return (top > 0) ? int(NavInd(top + 1)) : -1;
0604   }
0605
0606   VECGEOM_FORCE_INLINE
0607   VECCORE_ATT_HOST_DEVICE
0608   static void ReadTransformation(NavIndex_t nav_ind, Transformation3D &trans)
0609   {
0610     VECGEOM_VALIDATE(trans.IsIdentity(), << "ReadTransformation: destination must be an identity");
0611     auto record       = NavIndAddr(nav_ind);
0612     auto content_lhtr = reinterpret_cast<const unsigned char *>(record + 6);
0613     bool has_trans    = *(content_lhtr + 2) > 0;
0614     bool has_rot      = *(content_lhtr + 3) > 0;
0615     if (!(has_trans | has_rot)) return; // identity
0616     auto offset  = *(content_lhtr + 1);
0617     auto address = reinterpret_cast<const Precision *>(record + offset);
0618     VECGEOM_ASSERT(reinterpret_cast<uintptr_t>(address) % sizeof(Precision) == 0 &&
0619                    "ReadTransformation: transformation storage not aligned");
0620     trans.Set(address, address + int{has_trans} * 3, has_trans, has_rot);
0621   }
0622
0623   // FIXME THIS WILL BE A PROBLEM BECAUSE OF THE PRECISION REINTERPRET CAST
0624   template <typename Real_t>
0625   VECGEOM_FORCE_INLINE VECCORE_ATT_HOST_DEVICE static void ReadTransformation(NavIndex_t nav_ind,
0626                                                                               Transformation3DMP<Real_t> &trans)
0627   {
0628     VECGEOM_VALIDATE(trans.IsIdentity(), << "ReadTransformation: destination must be an identity");
0629     auto record       = NavIndAddr(nav_ind);
0630     auto content_lhtr = reinterpret_cast<const unsigned char *>(record + 6);
0631     bool has_trans    = *(content_lhtr + 2) > 0;
0632     bool has_rot      = *(content_lhtr + 3) > 0;
0633     if (!(has_trans | has_rot)) return; // identity
0634     auto offset  = *(content_lhtr + 1);
0635     auto address = reinterpret_cast<const Precision *>(record + offset);
0636     VECGEOM_ASSERT(reinterpret_cast<uintptr_t>(address) % sizeof(Precision) == 0 &&
0637                    "ReadTransformation: transformation storage not aligned");
0638     trans.Set(address, address + int{has_trans} * 3, has_trans, has_rot);
0639   }
0640
0641   VECCORE_ATT_HOST_DEVICE
0642   static void TopMatrixImpl(NavTuple_t const &nav_tuple, Transformation3D &trans)
0643   {
0644     // Get the multiplication of all parent scenes transformations, then multiply with the current
0645     // m_0 * m_1 * ... * m_n
0646     auto top = nav_tuple.Top();
0647     if (!top) return;
0648     ReadTransformation(top, trans);
0649     for (int ituple = nav_tuple.fLevel - 1; ituple >= 0; --ituple) {
0650       Transformation3D scene_trans;
0651       top = nav_tuple[ituple];
0652       ReadTransformation(top, scene_trans);
0653       trans *= scene_trans;
0654     }
0655   }
0656
0657   template <typename Real_t>
0658   VECCORE_ATT_HOST_DEVICE static void TopMatrixImpl(NavTuple_t const &nav_tuple, Transformation3DMP<Real_t> &trans)
0659   {
0660     // Get the multiplication of all parent scenes transformations, then multiply with the current
0661     // m_0 * m_1 * ... * m_n
0662     auto top = nav_tuple.Top();
0663     if (!top) return;
0664     ReadTransformation(top, trans);
0665     for (int ituple = nav_tuple.fLevel - 1; ituple >= 0; --ituple) {
0666       Transformation3DMP<Real_t> scene_trans;
0667       top = nav_tuple[ituple];
0668       ReadTransformation(top, scene_trans);
0669       trans *= scene_trans;
0670     }
0671   }
0672
0673   template <typename Real_t>
0674   VECCORE_ATT_HOST_DEVICE static void TopInSceneMatrixImpl(NavTuple_t const &nav_tuple,
0675                                                            Transformation3DMP<Real_t> &trans)
0676   {
0677     // Get transformation of the node in the top scene
0678     auto top = nav_tuple.Top();
0679     if (!top) return;
0680     ReadTransformation(top, trans);
0681   }
0682
0683   VECCORE_ATT_HOST_DEVICE
0684   static void TopInSceneMatrixImpl(NavTuple_t const &nav_tuple, Transformation3D &trans)
0685   {
0686     // Get transformation of the node in the top scene
0687     auto top = nav_tuple.Top();
0688     if (!top) return;
0689     ReadTransformation(top, trans);
0690   }
0691
0692   VECCORE_ATT_HOST_DEVICE
0693   static void SceneMatrixImpl(NavTuple_t const &nav_tuple, Transformation3D &trans)
0694   {
0695     // Get the top scene transformation
0696     auto top = nav_tuple.Top();
0697     if (!top) return;
0698     for (int ituple = nav_tuple.fLevel - 1; ituple >= 0; --ituple) {
0699       Transformation3D scene_trans;
0700       top = nav_tuple[ituple];
0701       ReadTransformation(top, scene_trans);
0702       trans *= scene_trans;
0703     }
0704   }
0705
0706   template <typename Real_t>
0707   VECCORE_ATT_HOST_DEVICE static void SceneMatrixImpl(NavTuple_t const &nav_tuple, Transformation3DMP<Real_t> &trans)
0708   {
0709     // Get the top scene transformation
0710     auto top = nav_tuple.Top();
0711     if (!top) return;
0712     for (int ituple = nav_tuple.fLevel - 1; ituple >= 0; --ituple) {
0713       Transformation3DMP<Real_t> scene_trans;
0714       top = nav_tuple[ituple];
0715       ReadTransformation(top, scene_trans);
0716       trans *= scene_trans;
0717     }
0718   }
0719
0720   VECCORE_ATT_HOST_DEVICE
0721   static Vector3D<Precision> GlobalToLocalImpl(NavTuple_t const &nav_tuple, Vector3D<Precision> const &globalpoint)
0722   {
0723     Transformation3D trans;
0724     TopMatrixImpl(nav_tuple, trans);
0725     Vector3D<Precision> local = trans.Transform(globalpoint);
0726     return local;
0727   }
0728
0729   template <typename Real_t>
0730   VECCORE_ATT_HOST_DEVICE static Vector3D<Real_t> GlobalToLocalImpl(NavTuple_t const &nav_tuple,
0731                                                                     Vector3D<Real_t> const &globalpoint)
0732   {
0733     Transformation3DMP<Real_t> trans;
0734     TopMatrixImpl(nav_tuple, trans);
0735     Vector3D<Real_t> local = trans.Transform(globalpoint);
0736     return local;
0737   }
0738
0739   // Intrerface methods
0740   VECGEOM_FORCE_INLINE
0741   VECCORE_ATT_HOST_DEVICE
0742   VPlacedVolume const *GetLastExited() const { return TopImpl(fLastExited); }
0743
0744   VECGEOM_FORCE_INLINE
0745   VECCORE_ATT_HOST_DEVICE
0746   NavTuple_t GetLastExitedState() const { return fLastExited; }
0747
0748   VECGEOM_FORCE_INLINE
0749   VECCORE_ATT_HOST_DEVICE
0750   int GetLastIdExited() const { return TopIdImpl(fLastExited); }
0751
0752   VECGEOM_FORCE_INLINE
0753   VECCORE_ATT_HOST_DEVICE
0754   void SetLastExited() { fLastExited = fNavTuple; }
0755
0756   VECGEOM_FORCE_INLINE
0757   VECCORE_ATT_HOST_DEVICE
0758   void SetLastExited(NavTuple_t const &nav_tuple) { fLastExited = nav_tuple; }
0759
0760   VECGEOM_FORCE_INLINE
0761   VECCORE_ATT_HOST_DEVICE
0762   void SetNavIndex(NavTuple_t const &nav_tuple) { fNavTuple = nav_tuple; }
0763
0764   VECGEOM_FORCE_INLINE
0765   VECCORE_ATT_HOST_DEVICE
0766   void SetNavIndex(NavIndex_t const &nav_index) { fNavTuple.Set(nav_index); }
0767
0768   VECCORE_ATT_HOST_DEVICE
0769   VECGEOM_FORCE_INLINE
0770   unsigned int GetLogicalId() const { return GetLogicalIdImpl(fNavTuple); }
0771
0772   VECCORE_ATT_HOST_DEVICE
0773   VECGEOM_FORCE_INLINE
0774   int GetChildId() const { return GetChildIdImpl(fNavTuple); }
0775
0776   VECGEOM_FORCE_INLINE
0777   VECCORE_ATT_HOST_DEVICE
0778   bool IsScene() const { return IsSceneImpl(fNavTuple); }
0779
0780   VECGEOM_FORCE_INLINE
0781   VECCORE_ATT_HOST_DEVICE
0782   bool IsDescendent(NavTuple_t const &parent) const { return IsDescendentImpl(fNavTuple, parent); }
0783
0784   VECCORE_ATT_HOST_DEVICE
0785   VECGEOM_FORCE_INLINE
0786   unsigned int GetNdaughters() const { return GetNdaughtersImpl(fNavTuple); }
0787
0788   VECGEOM_FORCE_INLINE
0789   VECCORE_ATT_HOST_DEVICE
0790   NavIndex_t GetId() const { return GetIdImpl(fNavTuple.Top()); }
0791
0792   VECGEOM_FORCE_INLINE
0793   VECCORE_ATT_HOST_DEVICE
0794   NavIndex_t GetParentSceneTopId() const
0795   {
0796     return (fNavTuple.fLevel > 0) ? GetIdImpl(fNavTuple[fNavTuple.fLevel - 1]) : 0;
0797   }
0798
0799   VECCORE_ATT_HOST_DEVICE
0800   VECGEOM_FORCE_INLINE
0801   bool GetSceneId(unsigned short &scene_id, unsigned short &newscene_id) const
0802   {
0803     return GetSceneIdImpl(fNavTuple, scene_id, newscene_id);
0804   }
0805
0806   VECCORE_ATT_HOST_DEVICE
0807   VECGEOM_FORCE_INLINE
0808   unsigned int GetSceneLevel() const { return GetSceneLevelImpl(fNavTuple); }
0809
0810   VECCORE_ATT_HOST_DEVICE
0811   VECGEOM_FORCE_INLINE
0812   unsigned short GetParentScene() const { return GetParentSceneImpl(fNavTuple); }
0813
0814   VECGEOM_FORCE_INLINE
0815   VECCORE_ATT_HOST_DEVICE
0816   void Push(VPlacedVolume const *v) { PushImpl(fNavTuple, v); }
0817
0818   VECGEOM_FORCE_INLINE
0819   VECCORE_ATT_HOST_DEVICE
0820   void Push(int iplaced) { PushImpl(fNavTuple, iplaced); }
0821
0822   VECGEOM_FORCE_INLINE
0823   VECCORE_ATT_HOST_DEVICE
0824   void PushDaughter(int idaughter) { PushDaughterImpl(fNavTuple, idaughter); }
0825
0826   VECGEOM_FORCE_INLINE
0827   VECCORE_ATT_HOST_DEVICE
0828   void PushScene(NavIndex_t nav_ind)
0829   {
0830     fNavTuple.fLevel++;
0831     fNavTuple.Set(nav_ind);
0832   }
0833
0834   VECGEOM_FORCE_INLINE
0835   VECCORE_ATT_HOST_DEVICE
0836   void Pop() { PopImpl(fNavTuple); }
0837
0838   VECGEOM_FORCE_INLINE
0839   VECCORE_ATT_HOST_DEVICE
0840   void PopScene()
0841   {
0842     if (fNavTuple.fLevel > 0) fNavTuple.fLevel--;
0843   }
0844
0845   VECGEOM_FORCE_INLINE
0846   VECCORE_ATT_HOST_DEVICE
0847   VPlacedVolume const *Top() const { return TopImpl(fNavTuple); }
0848
0849   VECGEOM_FORCE_INLINE
0850   VECCORE_ATT_HOST_DEVICE
0851   int TopId() const { return TopIdImpl(fNavTuple); }
0852
0853   /**
0854    * returns the number of FILLED LEVELS such that
0855    * state.GetNode( state.GetLevel() ) == state.Top()
0856    */
0857   VECGEOM_FORCE_INLINE
0858   VECCORE_ATT_HOST_DEVICE
0859   unsigned char GetLevel() const { return GetLevelImpl(fNavTuple); }
0860
0861   /** Compatibility getter for NavigationState interface */
0862   VECGEOM_FORCE_INLINE
0863   VECCORE_ATT_HOST_DEVICE
0864   unsigned char GetCurrentLevel() const { return GetLevel() + 1; }
0865
0866   /**
0867    * Returns the navigation index for a level smaller/equal than the current level.
0868    */
0869   VECCORE_ATT_HOST_DEVICE
0870   VECGEOM_FORCE_INLINE
0871   NavTuple_t GetNavTuple(int level) const { return GetNavTupleImpl(fNavTuple, level); }
0872
0873   /**
0874    * Returns the placed volume at a evel smaller/equal than the current level.
0875    */
0876   VECGEOM_FORCE_INLINE
0877   VECCORE_ATT_HOST_DEVICE
0878   VPlacedVolume const *At(int level) const
0879   {
0880     auto nav_ind = GetNavTupleImpl(fNavTuple, level).Top();
0881     return (nav_ind > 0) ? ToPlacedVolume(NavInd(nav_ind + 1)) : nullptr;
0882   }
0883
0884   /**
0885    * Returns the index of a placed volume at a evel smaller/equal than the current level.
0886    */
0887   VECGEOM_FORCE_INLINE
0888   VECCORE_ATT_HOST_DEVICE
0889   size_t ValueAt(int level) const
0890   {
0891     auto nav_ind = GetNavTupleImpl(fNavTuple, level).Top();
0892     return (nav_ind > 0) ? size_t(NavInd(nav_ind + 1)) : 0;
0893   }
0894
0895   VECCORE_ATT_HOST_DEVICE
0896   void TopMatrix(Transformation3D &trans) const { TopMatrixImpl(fNavTuple, trans); }
0897
0898   template <typename Real_t>
0899   VECCORE_ATT_HOST_DEVICE void TopMatrix(Transformation3DMP<Real_t> &trans) const
0900   {
0901     TopMatrixImpl(fNavTuple, trans);
0902   }
0903
0904   VECCORE_ATT_HOST_DEVICE
0905   void TopInSceneMatrix(Transformation3D &trans) const { TopInSceneMatrixImpl(fNavTuple, trans); }
0906
0907   template <typename Real_t>
0908   VECCORE_ATT_HOST_DEVICE void TopInSceneMatrix(Transformation3DMP<Real_t> &trans) const
0909   {
0910     TopInSceneMatrixImpl(fNavTuple, trans);
0911   }
0912
0913   VECCORE_ATT_HOST_DEVICE
0914   void SceneMatrix(Transformation3D &trans) const { SceneMatrixImpl(fNavTuple, trans); }
0915
0916   template <typename Real_t>
0917   VECCORE_ATT_HOST_DEVICE void SceneMatrix(Transformation3DMP<Real_t> &trans) const
0918   {
0919     SceneMatrixImpl(fNavTuple, trans);
0920   }
0921
0922   VECCORE_ATT_HOST_DEVICE
0923   void TopMatrix(int tolevel, Transformation3D &trans) const
0924   {
0925     TopMatrixImpl(GetNavTupleImpl(fNavTuple, tolevel), trans);
0926   }
0927
0928   template <typename Real_t>
0929   VECCORE_ATT_HOST_DEVICE void TopMatrix(int tolevel, Transformation3DMP<Real_t> &trans) const
0930   {
0931     TopMatrixImpl(GetNavTupleImpl(fNavTuple, tolevel), trans);
0932   }
0933
0934   // returning a "delta" transformation that can transform
0935   // coordinates given in reference frame of this->Top() to the reference frame of other->Top()
0936   // simply with otherlocalcoordinate = delta.Transform( thislocalcoordinate )
0937   VECCORE_ATT_HOST_DEVICE
0938   void DeltaTransformation(NavStateTuple const &other, Transformation3D &delta) const
0939   {
0940     Transformation3D g2;
0941     Transformation3D g1;
0942     other.TopMatrix(g2);
0943     this->TopMatrix(g1);
0944     delta = g1.Inverse();
0945     // Trans/rot properties already correctly set
0946     // g2.SetProperties();
0947     // delta.SetProperties();
0948     delta.FixZeroes();
0949     delta.MultiplyFromRight(g2);
0950     delta.FixZeroes();
0951   }
0952
0953   template <typename Real_t>
0954   VECCORE_ATT_HOST_DEVICE void DeltaTransformation(NavStateTuple const &other, Transformation3DMP<Real_t> &delta) const
0955   {
0956     Transformation3DMP<Real_t> g2;
0957     Transformation3DMP<Real_t> g1;
0958     other.TopMatrix(g2);
0959     this->TopMatrix(g1);
0960     delta = g1.Inverse();
0961     // Trans/rot properties already correctly set
0962     // g2.SetProperties();
0963     // delta.SetProperties();
0964     delta.FixZeroes();
0965     delta.MultiplyFromRight(g2);
0966     delta.FixZeroes();
0967   }
0968
0969   // VECGEOM_FORCE_INLINE
0970   VECCORE_ATT_HOST_DEVICE
0971   Vector3D<Precision> GlobalToLocal(Vector3D<Precision> const &localpoint) const
0972   {
0973     return GlobalToLocalImpl(fNavTuple, localpoint);
0974   }
0975
0976   VECCORE_ATT_HOST_DEVICE
0977   Vector3D<Precision> GlobalToLocal(Vector3D<Precision> const &localpoint, int tolevel) const
0978   {
0979     return GlobalToLocalImpl(GetNavTupleImpl(fNavTuple, tolevel), localpoint);
0980   }
0981
0982   /**
0983    * calculates if other navigation state takes a different branch in geometry path or is on same branch
0984    * ( two states are on same branch if one can connect the states just by going upwards or downwards ( or do nothing
0985    * ))
0986    */
0987   VECCORE_ATT_HOST_DEVICE
0988   int Distance(NavStateTuple const &other) const
0989   {
0990     int lastcommonlevel = -1;
0991     int thislevel       = GetLevel();
0992     int otherlevel      = other.GetLevel();
0993     int maxlevel        = Min(thislevel, otherlevel);
0994
0995     //  algorithm: start on top and go down until paths split
0996     for (int i = 0; i < maxlevel + 1; i++) {
0997       if (this->At(i) == other.At(i)) {
0998         lastcommonlevel = i;
0999       } else {
1000         break;
1001       }
1002     }
1003
1004     return (thislevel - lastcommonlevel) + (otherlevel - lastcommonlevel);
1005   }
1006
1007   // returns a string representation of a (relative) sequence of operations/moves
1008   // that transforms this navigation state into the other navigation state
1009   // example:
1010   // state1 = /0/1/1/
1011   // state2 = /0/2/2/3
1012   // results in string
1013   // "/up/horiz/1/down/2/down/3" with 4 operations "up", "horiz", "down", "down"
1014   // the sequence of moves is the following
1015   // up: /0/1/1 --> /0/1/
1016   // horiz/1 : 0/1 --> /0/2 ( == /0/(1+1) )   "we are hopping from daughter 1 to 2 (which corresponds to a step of 1)"
1017   // down/2 : /0/2 --> /0/2/2   "going further down 2nd daughter"
1018   // down/3 : /0/2/2/3 --> /0/2/2/3  "going further down 2nd daughter"
1019   std::string RelativePath(NavStateTuple const &other) const
1020   {
1021     int lastcommonlevel = -1;
1022     int thislevel       = GetLevel();
1023     int otherlevel      = other.GetLevel();
1024     int maxlevel        = Min(thislevel, otherlevel);
1025     std::stringstream str;
1026     //  algorithm: start on top and go down until paths split
1027     for (int i = 0; i < maxlevel + 1; i++) {
1028       if (this->At(i) == other.At(i)) {
1029         lastcommonlevel = i;
1030       } else {
1031         break;
1032       }
1033     }
1034
1035     // paths are the same
1036     if (thislevel == lastcommonlevel && otherlevel == lastcommonlevel) {
1037       return std::string("");
1038     }
1039
1040     // emit only ups
1041     if (thislevel > lastcommonlevel && otherlevel == lastcommonlevel) {
1042       for (int i = 0; i < thislevel - lastcommonlevel; ++i) {
1043         str << "/up";
1044       }
1045       return str.str();
1046     }
1047
1048     // emit only downs
1049     if (thislevel == lastcommonlevel && otherlevel > lastcommonlevel) {
1050       for (int i = lastcommonlevel + 1; i <= otherlevel; ++i) {
1051         str << "/down";
1052         str << "/" << other.ValueAt(i);
1053       }
1054       return str.str();
1055     }
1056
1057     // mixed case: first up; then down
1058     if (thislevel > lastcommonlevel && otherlevel > lastcommonlevel) {
1059       // emit ups
1060       int level = thislevel;
1061       for (; level > lastcommonlevel + 1; --level) {
1062         str << "/up";
1063       }
1064
1065       level = lastcommonlevel + 1;
1066       // emit horiz ( exists when there is a turning point )
1067       int delta = other.ValueAt(level) - this->ValueAt(level);
1068       if (delta != 0) str << "/horiz/" << delta;
1069
1070       level++;
1071       // emit downs with index
1072       for (; level <= otherlevel; ++level) {
1073         str << "/down/" << other.ValueAt(level);
1074       }
1075     }
1076     return str.str();
1077   }
1078
1079   // functions useful to "serialize" navigationstate
1080   // the Vector-of-Indices basically describes the path on the tree taken from top to bottom
1081   // an index corresponds to a daughter
1082
1083   void GetPathAsListOfIndices(std::list<uint> &indices) const
1084   {
1085     indices.clear();
1086     if (IsOutside()) return;
1087
1088     auto nav_tuple = fNavTuple;
1089     while (nav_tuple.Top() > 1) {
1090       auto pvol = TopImpl(nav_tuple);
1091       indices.push_front(pvol->GetChildId());
1092       PopImpl(nav_tuple);
1093     }
1094     // Paths start always with 0
1095     indices.push_front(0);
1096   }
1097
1098   void ResetPathFromListOfIndices(VPlacedVolume const *world, std::list<uint> const &indices)
1099   {
1100     // clear current nav state
1101     Clear();
1102     auto vol    = world;
1103     int counter = 0;
1104     for (auto id : indices) {
1105       if (counter > 0) vol = vol->GetDaughters().operator[](id);
1106       Push(vol);
1107       counter++;
1108     }
1109   }
1110
1111   // replaces the volume pointers from CPU volumes in fPath
1112   // to the equivalent pointers on the GPU
1113   // uses the CudaManager to do so
1114   void ConvertToGPUPointers() {}
1115
1116   // replaces the pointers from GPU volumes in fPath
1117   // to the equivalent pointers on the CPU
1118   // uses the CudaManager to do so
1119   void ConvertToCPUPointers() {}
1120
1121   // clear all information
1122   VECGEOM_FORCE_INLINE
1123   VECCORE_ATT_HOST_DEVICE
1124   void Clear()
1125   {
1126     fNavTuple.Clear();
1127     fLastExited.Clear();
1128     fOnBoundary = false;
1129   }
1130
1131   VECCORE_ATT_HOST_DEVICE
1132   static void PrintRecord(NavIndex_t nav_ind)
1133   {
1134     if (nav_ind == 0) return;
1135     auto parent       = NavInd(nav_ind);
1136     auto placed_id    = NavInd(nav_ind + 1);
1137     auto child_id     = NavInd(nav_ind + 2);
1138     auto id           = NavInd(nav_ind + 3);
1139     auto logical_addr = NavInd(nav_ind + 4);
1140     auto logical_id   = NavInd(logical_addr);
1141     auto scenes       = reinterpret_cast<const unsigned short *>(NavIndAddr(nav_ind + 5));
1142     auto scene_id     = scenes[0];
1143     auto newscene_id  = scenes[1];
1144     auto level        = GetLevelImpl(nav_ind);
1145     auto nd           = NavInd(logical_addr + 1);
1146     printf("| navind %u |+0| parent %u |+1| placed_id %u |+2| child_id %u |+3| id %u |+4| logical_addr %u |+5| scene "
1147            "%hu | "
1148            "new_scene %hu |+6| level %u | ... |%u| logical_id %u |+1| nd %u ",
1149            nav_ind, parent, placed_id, child_id, id, logical_addr, scene_id, newscene_id, level, logical_addr,
1150            logical_id, nd);
1151     if (nd > 0) printf("|+2| d0 %u | ...", NavInd(logical_addr + 2));
1152     printf("\n");
1153   }
1154
1155   VECCORE_ATT_HOST_DEVICE
1156   void Print(bool print_names = false) const
1157   {
1158     if (fNavTuple.Top() == 0 && fNavTuple.fLevel == 0) {
1159       printf("navInd=0, id=0, path=outside\n");
1160       return;
1161     }
1162     NavTuple_t nav_tuple;
1163     auto level = GetLevel();
1164     printf("navInd=");
1165     for (unsigned i = 0; i <= fNavTuple.fLevel; ++i) {
1166       unsigned short scene_id = 0, newscene_id = 0;
1167       nav_tuple.Push(fNavTuple[i]);
1168       GetSceneIdImpl(nav_tuple, scene_id, newscene_id);
1169       printf("s%u:%u", scene_id, fNavTuple[i]);
1170       if (i < fNavTuple.fLevel) printf(" | ");
1171     }
1172     printf(" lastExited=");
1173     for (unsigned i = 0; i <= fLastExited.fLevel; ++i) {
1174       printf("%u", fLastExited[i]);
1175       if (i < fLastExited.fLevel) printf(" | ");
1176     }
1177     printf(", id=%u, level=%u/%u,  onBoundary=%s, path=<", GetId(), level, GetMaxLevel(),
1178            (fOnBoundary ? "true" : "false"));
1179     int last_scene = 0;
1180     nav_tuple.Clear();
1181     for (int i = 0; i <= level; ++i) {
1182       unsigned short scene_id = 0, newscene_id = 0;
1183       nav_tuple = GetNavTupleImpl(fNavTuple, i);
1184       GetSceneIdImpl(nav_tuple, scene_id, newscene_id);
1185       if (scene_id != last_scene) {
1186         printf(" | ");
1187         last_scene = scene_id;
1188       }
1189 #ifndef VECCORE_CUDA
1190       if (print_names) {
1191         auto vol = At(i);
1192         printf("/%s", vol ? vol->GetLabel().c_str() : "TOP_SCENE");
1193       } else
1194 #endif
1195         printf("/%u", nav_tuple.Top());
1196     }
1197     printf(">\n");
1198   }
1199
1200   VECCORE_ATT_HOST_DEVICE
1201   static void PrintTopImpl(NavIndex_t top)
1202   {
1203     if (top == 0) {
1204       printf("navInd=0, id=0, path=TOP_SCENE\n");
1205       return;
1206     }
1207     // Find level
1208     int level         = 0;
1209     NavIndex_t mother = top;
1210     while (mother) {
1211       level++;
1212       mother = NavInd(mother);
1213     }
1214     printf("navInd=");
1215     unsigned short scene_id = 0, newscene_id = 0;
1216     GetSceneIdImpl(top, scene_id, newscene_id);
1217     printf("s%u:%u", scene_id, top);
1218
1219     printf(", id=%u, level=%u,  path=<", GetIdImpl(top), level);
1220     for (int i = 0; i <= level; ++i) {
1221       mother = top;
1222       for (int j = 0; j < level - i; ++j)
1223         mother = NavInd(mother);
1224 #ifndef VECCORE_CUDA
1225       auto vol = (mother > 0) ? ToPlacedVolume(NavInd(mother + 1)) : nullptr;
1226       printf("/%s", vol ? vol->GetLabel().c_str() : "TOP_SCENE");
1227 #else
1228       printf("/%u", mother);
1229 #endif
1230     }
1231     printf(">\n");
1232   }
1233
1234   VECCORE_ATT_HOST_DEVICE
1235   void PrintTop() const
1236   {
1237     if (fNavTuple.Top() == 0) {
1238       printf("navInd=0, id=0, path=outside\n");
1239       return;
1240     }
1241     auto level = GetLevel();
1242     printf("navInd=");
1243     auto top                = fNavTuple.Top();
1244     unsigned short scene_id = 0, newscene_id = 0;
1245     GetSceneIdImpl(fNavTuple, scene_id, newscene_id);
1246     printf("s%u:%u", scene_id, top);
1247
1248     printf(", id=%u, level=%u/%u,  onBoundary=%s, path=<", GetId(), level, GetMaxLevel(),
1249            (fOnBoundary ? "true" : "false"));
1250     int top_scene = scene_id;
1251     for (int i = 0; i <= level; ++i) {
1252       auto nav_tuple = GetNavTupleImpl(fNavTuple, i);
1253       GetSceneIdImpl(nav_tuple, scene_id, newscene_id);
1254       if (scene_id != top_scene) continue;
1255 #ifndef VECCORE_CUDA
1256       auto vol = At(i);
1257       printf("/%s", vol ? vol->GetLabel().c_str() : "TOP_SCENE");
1258 #else
1259       printf("/%u", nav_tuple.Top());
1260 #endif
1261     }
1262     printf(">\n");
1263   }
1264
1265   VECCORE_ATT_HOST_DEVICE
1266   void Dump() const { Print(); }
1267
1268   VECGEOM_FORCE_INLINE
1269   VECCORE_ATT_HOST_DEVICE
1270   bool HasSamePathAsOther(NavStateTuple const &other) const { return (fNavTuple == other.fNavTuple); }
1271
1272   void printValueSequence(std::ostream & = std::cerr) const;
1273
1274   // calculates a checksum along the path
1275   // can be used (as a quick criterion) to see whether 2 states are same
1276   unsigned long getCheckSum() const
1277   {
1278     unsigned long checksum = 0;
1279     for (uint ituple = 0; ituple <= fNavTuple.fLevel; ++ituple)
1280       checksum += (unsigned long)fNavTuple[ituple];
1281
1282     return checksum;
1283   }
1284
1285   /**
1286     function returning whether the point (current navigation state) is outside the detector setup
1287   */
1288   VECGEOM_FORCE_INLINE
1289   VECCORE_ATT_HOST_DEVICE
1290   bool IsOutside() const { return (fNavTuple.Top() == 0); }
1291
1292   VECGEOM_FORCE_INLINE
1293   VECCORE_ATT_HOST_DEVICE
1294   bool IsOnBoundary() const { return fOnBoundary; }
1295
1296   VECGEOM_FORCE_INLINE
1297   VECCORE_ATT_HOST_DEVICE
1298   void SetBoundaryState(bool b) { fOnBoundary = b; }
1299 };
1300
1301 /**
1302  * encodes the geometry path as a concatenated string of ( Value_t ) present in fPath
1303  */
1304 inline void NavStateTuple::printValueSequence(std::ostream &stream) const
1305 {
1306   auto level = GetLevel();
1307   for (int i = 0; i < level + 1; ++i) {
1308     auto pvol = At(i);
1309     if (pvol) stream << "/" << ValueAt(i) << "(" << pvol->GetLabel() << ")";
1310   }
1311 }
1312
1313 } // namespace vecgeom
1314
1315 #endif // VECGEOM_NAVIGATION_NAVSTATETUPLE_H_