VecGeom/navigation/NavStatePath.h

0001 /// \file NavStatePath.h
0002 /// \author Sandro Wenzel (sandro.wenzel@cern.ch)
0003 /// \date 12.03.2014
0004
0005 #ifndef VECGEOM_NAVIGATION_NAVSTATEPATH_H_
0006 #define VECGEOM_NAVIGATION_NAVSTATEPATH_H_
0007
0008 #include "VecGeom/base/Config.h"
0009
0010 #include "VecGeom/base/VariableSizeObj.h"
0011 #include "VecGeom/base/Transformation3D.h"
0012 #include "VecGeom/volumes/PlacedVolume.h"
0013 #include "VecGeom/management/GeoManager.h"
0014 #ifdef VECGEOM_ENABLE_CUDA
0015 #include "VecGeom/management/CudaManager.h"
0016 #endif
0017 #include "VecGeom/base/Global.h"
0018
0019 #include <iostream>
0020 #include <string>
0021
0022 class TGeoBranchArray;
0023
0024 // gcc 4.8.2's -Wnon-virtual-dtor is broken and turned on by -Weffc++, we
0025 // need to disable it for SOA3D
0026
0027 #if __GNUC__ < 3 || (__GNUC__ == 4 && __GNUC_MINOR__ <= 8)
0028
0029 #pragma GCC diagnostic push
0030 #pragma GCC diagnostic ignored "-Wnon-virtual-dtor"
0031 #pragma GCC diagnostic ignored "-Weffc++"
0032 #define GCC_DIAG_POP_NEEDED
0033 #endif
0034
0035 namespace vecgeom {
0036 inline namespace VECGEOM_IMPL_NAMESPACE {
0037
0038 // the NavStateIndex type determines is used
0039 // to calculate addresses of PlacedVolumes
0040 // a short type should be used in case the number of PlacedVolumes can
0041 // be counted with 16bits
0042 // TODO: consider putting uint16 + uint32 types
0043 #ifdef VECGEOM_USE_INDEXEDNAVSTATES
0044 // typedef unsigned short NavStateIndex_t;
0045 typedef unsigned long NavStateIndex_t;
0046 #else
0047 typedef VPlacedVolume const *NavStateIndex_t;
0048 #endif
0049
0050 // helper functionality to convert from NavStateIndex_t to *PlacedVolumes and back
0051 // the template abstraction also allows to go back to pointers as NavStateIndex_t
0052 // via a template specialization
0053 // T stands for NavStateIndex_t
0054 template <typename T>
0055 struct Index2PVolumeConverter {
0056   VECCORE_ATT_HOST_DEVICE
0057   VECGEOM_FORCE_INLINE
0058   static VPlacedVolume const *ToPlacedVolume(T index)
0059   {
0060 #ifdef VECCORE_CUDA_DEVICE_COMPILATION
0061     // checking here for NVCC_DEVICE since the global variable globaldevicegeomgata::gCompact...
0062     // is marked __device__ and can only be compiled within device compiler passes
0063     VECGEOM_ASSERT(vecgeom::globaldevicegeomdata::gCompactPlacedVolBuffer != nullptr);
0064     return &vecgeom::globaldevicegeomdata::gCompactPlacedVolBuffer[index];
0065 #else
0066 #ifndef VECCORE_CUDA
0067     VECGEOM_ASSERT(vecgeom::GeoManager::gCompactPlacedVolBuffer == nullptr ||
0068                    vecgeom::GeoManager::gCompactPlacedVolBuffer[index].id() == index);
0069     return &vecgeom::GeoManager::gCompactPlacedVolBuffer[index];
0070 #else
0071     // this is the case when we compile with nvcc for host side
0072     // (failed previously due to undefined symbol vecgeom::cuda::GeoManager::gCompactPlacedVolBuffer)
0073     VECGEOM_VALIDATE(false, << "reached unimplement code");
0074     (void)index; // avoid unused parameter warning.
0075     return nullptr;
0076 #endif
0077 #endif
0078   }
0079
0080   VECCORE_ATT_HOST_DEVICE
0081   VECGEOM_FORCE_INLINE
0082   static T ToIndex(VPlacedVolume const *pvol) { return pvol->id(); }
0083 };
0084
0085 // template specialization when we directly save VPlacedVolume pointers into the NavStates
0086 template <>
0087 struct Index2PVolumeConverter<VPlacedVolume const *> {
0088   VECCORE_ATT_HOST_DEVICE
0089   VECGEOM_FORCE_INLINE
0090   static VPlacedVolume const *ToPlacedVolume(VPlacedVolume const *pvol) { return pvol; }
0091   VECCORE_ATT_HOST_DEVICE
0092   VECGEOM_FORCE_INLINE
0093   static VPlacedVolume const *ToIndex(VPlacedVolume const *pvol) { return pvol; }
0094 };
0095
0096 /**
0097  * A class describing a current geometry state
0098  * likely there will be such an object for each particle/track currently treated.
0099  */
0100 class NavStatePath : protected VariableSizeObjectInterface<NavStatePath, NavStateIndex_t>,
0101                      private Index2PVolumeConverter<NavStateIndex_t> {
0102 public:
0103   using Value_t        = NavStateIndex_t;
0104   using Base_t         = VariableSizeObjectInterface<NavStatePath, Value_t>;
0105   using VariableData_t = VariableSizeObj<Value_t>;
0106
0107 private:
0108   friend Base_t;
0109
0110   // Required by VariableSizeObjectInterface
0111   VECCORE_ATT_HOST_DEVICE
0112   VariableData_t &GetVariableData() { return fPath; }
0113   VECCORE_ATT_HOST_DEVICE
0114   const VariableData_t &GetVariableData() const { return fPath; }
0115
0116   unsigned char
0117       fCurrentLevel; // value indicating the next free slot in the fPath array ( ergo the current geometry depth )
0118   // we choose unsigned char in order to save memory, thus supporting geometry depths up to 255 which seems large enough
0119
0120   // a member to cache some state information across state usages
0121   // one particular example could be a calculated index for this state
0122   // if fCache == -1 it means the we have to recalculate it; otherwise we don't
0123   short fCache;
0124
0125   bool fOnBoundary; // flag indicating whether track is on boundary of the "Top()" placed volume
0126
0127 #ifdef VECGEOM_USE_CACHED_TRANSFORMATIONS
0128   // Cached transformation for top level
0129   bool fCacheM = true; // flag indicating whether the global matrix for the state is cached
0130   Transformation3D fTopTrans;
0131 #endif
0132
0133   // pointer data follows; has to be last
0134   VariableSizeObj<Value_t> fPath;
0135
0136   // constructors and assignment operators are private
0137   // states have to be constructed using MakeInstance() function
0138   VECGEOM_FORCE_INLINE
0139   VECCORE_ATT_HOST_DEVICE
0140   NavStatePath(size_t nvalues);
0141
0142   VECGEOM_FORCE_INLINE
0143   VECCORE_ATT_HOST_DEVICE
0144   NavStatePath(size_t new_size, NavStatePath &other)
0145       : fCurrentLevel(other.fCurrentLevel), fCache(-1), fOnBoundary(other.fOnBoundary),
0146 #ifdef VECGEOM_USE_CACHED_TRANSFORMATIONS
0147         fCacheM(other.fCacheM), fTopTrans(other.fTopTrans),
0148 #endif
0149         fPath(new_size, other.fPath)
0150   {
0151     // Raw memcpy of the content to another existing state.
0152     //
0153     // in case NavStatePath was a virtual class: change to
0154     // std::memcpy(other->DataStart(), DataStart(), DataSize());
0155
0156     if (new_size > other.fPath.fN) {
0157       memset(fPath.GetValues() + other.fPath.fN, 0, new_size - other.fPath.fN);
0158     }
0159   }
0160
0161   // some private management methods
0162   VECGEOM_FORCE_INLINE
0163   VECCORE_ATT_HOST_DEVICE
0164   void InitInternalStorage();
0165
0166 private:
0167   // The data start should point to the address of the first data member,
0168   // after the virtual table
0169   // the purpose is probably for the Copy function
0170   const void *DataStart() const { return (const void *)&fCurrentLevel; }
0171   const void *ObjectStart() const { return (const void *)this; }
0172   void *DataStart() { return (void *)&fCurrentLevel; }
0173   void *ObjectStart() { return (void *)this; }
0174
0175   // The actual size of the data for an instance, excluding the virtual table
0176   size_t DataSize() const { return SizeOf() + (size_t)ObjectStart() - (size_t)DataStart(); }
0177
0178 public:
0179   // replaces the volume pointers from CPU volumes in fPath
0180   // to the equivalent pointers on the GPU
0181   // uses the CudaManager to do so
0182   void ConvertToGPUPointers();
0183
0184   // replaces the pointers from GPU volumes in fPath
0185   // to the equivalent pointers on the CPU
0186   // uses the CudaManager to do so
0187   void ConvertToCPUPointers();
0188
0189   // Enumerate the part of the private interface, we want to expose.
0190   using Base_t::MakeCopy;
0191   using Base_t::MakeCopyAt;
0192   using Base_t::ReleaseInstance;
0193   using Base_t::SizeOf;
0194   using Base_t::SizeOfAlignAware;
0195
0196   // Enumerate functions from converter which we want to use
0197   // ( without retyping of the struct name )
0198   using Index2PVolumeConverter<NavStateIndex_t>::ToIndex;
0199   using Index2PVolumeConverter<NavStateIndex_t>::ToPlacedVolume;
0200
0201   // produces a compact navigation state object of a certain depth
0202   // the caller can give a memory address where the object will
0203   // be placed
0204   // the caller has to make sure that the size of the external memory
0205   // is >= sizeof(NavStatePath) + sizeof(VPlacedVolume*)*maxlevel
0206   //
0207   // Methods MakeInstance(), MakeInstanceAt(), MakeCopy() and MakeCopyAt() are provided by
0208   // VariableSizeObjectInterface
0209
0210   VECCORE_ATT_HOST_DEVICE
0211   static NavStatePath *MakeInstance(int maxlevel)
0212   {
0213     // MaxLevel is 'zero' based (i.e. maxlevel==0 requires one value)
0214     return Base_t::MakeInstance(maxlevel + 1);
0215   }
0216
0217   VECCORE_ATT_HOST_DEVICE
0218   static NavStatePath *MakeInstanceAt(int maxlevel, void *addr)
0219   {
0220     // MaxLevel is 'zero' based (i.e. maxlevel==0 requires one value)
0221     return Base_t::MakeInstanceAt(maxlevel + 1, addr);
0222   }
0223
0224   // returns the size in bytes of a NavStatePath object with internal
0225   // path depth maxlevel
0226   VECCORE_ATT_HOST_DEVICE
0227   static size_t SizeOfInstance(int maxlevel)
0228   {
0229     // MaxLevel is 'zero' based (i.e. maxlevel==0 requires one value)
0230     return VariableSizeObjectInterface::SizeOf(maxlevel + 1);
0231   }
0232
0233   // returns the size in bytes of a NavStatePath object with internal
0234   // path depth maxlevel -- including space needed for padding to next aligned object
0235   // of same kind
0236   VECCORE_ATT_HOST_DEVICE
0237   static size_t SizeOfInstanceAlignAware(int maxlevel)
0238   {
0239     // MaxLevel is 'zero' based (i.e. maxlevel==0 requires one value)
0240     return VariableSizeObjectInterface::SizeOfAlignAware(maxlevel + 1);
0241   }
0242
0243   VECCORE_ATT_HOST_DEVICE
0244   int GetObjectSize() const { return SizeOf(GetMaxLevel()); }
0245
0246   VECCORE_ATT_HOST_DEVICE
0247   int SizeOf() const { return NavStatePath::SizeOfInstance(GetMaxLevel()); }
0248
0249   VECGEOM_FORCE_INLINE
0250   VECCORE_ATT_HOST_DEVICE
0251   NavStatePath &operator=(NavStatePath const &rhs);
0252
0253   // functions useful to "serialize" navigationstate
0254   // the Vector-of-Indices basically describes the path on the tree taken from top to bottom
0255   // an index corresponds to a daughter
0256   void GetPathAsListOfIndices(std::list<uint> &indices) const;
0257   void ResetPathFromListOfIndices(VPlacedVolume const *world, std::list<uint> const &indices);
0258
0259   VECCORE_ATT_HOST_DEVICE
0260   void CopyTo(NavStatePath *other) const
0261   {
0262     // Raw memcpy of the content to another existing state.
0263     //
0264     // in case NavStatePath was a virtual class: change to
0265     // std::memcpy(other->DataStart(), DataStart(), DataSize());
0266     bool alloc = other->fPath.fSelfAlloc;
0267     // std::memcpy(other, this, this->SizeOf());
0268     // we only need to copy to relevant depth
0269     // GetCurrentLevel indicates the 'next' level, i.e. currentLevel==0 is empty
0270     // fCurrentLevel = maxlevel+1 is full
0271     // SizeOfInstance expect [0,maxlevel] and add +1 to its params
0272     std::memcpy(static_cast<void *>(other), this, NavStatePath::SizeOfInstance(this->GetCurrentLevel()));
0273
0274     other->fPath.fSelfAlloc = alloc;
0275   }
0276
0277   // copies a fixed and predetermined number of bytes
0278   // might be useful for specialized navigators which know the depth + SizeOf in advance
0279   // N is number of bytes to be copied and can be obtained by a prior call to constexpr NavStatePath::SizeOf( ... );
0280   template <size_t N>
0281   void CopyToFixedSize(NavStatePath *other) const
0282   {
0283     bool alloc = other->fPath.fSelfAlloc;
0284     for (size_t i = 0; i < N; ++i) {
0285       ((char *)other)[i] = ((char *)this)[i];
0286     }
0287     other->fPath.fSelfAlloc = alloc;
0288   }
0289
0290   VECGEOM_FORCE_INLINE
0291   VECCORE_ATT_HOST_DEVICE
0292   ~NavStatePath();
0293
0294   // what else: operator new etc...
0295
0296   VECGEOM_FORCE_INLINE
0297   VECCORE_ATT_HOST_DEVICE
0298   unsigned char GetMaxLevel() const { return fPath.fN - 1; }
0299
0300   VECGEOM_FORCE_INLINE
0301   VECCORE_ATT_HOST_DEVICE
0302   unsigned char GetCurrentLevel() const { return fCurrentLevel; }
0303
0304   VECGEOM_FORCE_INLINE
0305   VECCORE_ATT_HOST_DEVICE
0306   VPlacedVolume const *GetLastExited() const
0307   { /*one beyond current*/
0308     // A.G this implementation was bugged, because it relies on the fact that this state object actually exited a deeper
0309     // path which may not be true - the path may be even set from outside. This implementation cannot provide this
0310     // functionality unless a specific data member is added and set via SetLastExited return (fCurrentLevel <
0311     // GetMaxLevel() - 1) ? ToPlacedVolume(fPath[fCurrentLevel]) : nullptr;
0312     return nullptr;
0313   }
0314
0315   VECGEOM_FORCE_INLINE
0316   VECCORE_ATT_HOST_DEVICE
0317   void SetLastExited() {}
0318
0319   // better to use pop and push
0320   VECGEOM_FORCE_INLINE
0321   VECCORE_ATT_HOST_DEVICE
0322   void Push(VPlacedVolume const *);
0323
0324   // Push a given daughter to the state
0325   VECGEOM_FORCE_INLINE
0326   VECCORE_ATT_HOST_DEVICE
0327   void Push(unsigned short);
0328
0329   // a push version operating on IndexTypes
0330   VECGEOM_FORCE_INLINE
0331   VECCORE_ATT_HOST_DEVICE
0332   void PushIndexType(NavStateIndex_t);
0333
0334   VECGEOM_FORCE_INLINE
0335   VECCORE_ATT_HOST_DEVICE
0336   VPlacedVolume const *Top() const;
0337
0338   VECGEOM_FORCE_INLINE
0339   VECCORE_ATT_HOST_DEVICE
0340   VPlacedVolume const *At(int level) const { return ToPlacedVolume(fPath[level]); }
0341
0342   VECGEOM_FORCE_INLINE
0343   VECCORE_ATT_HOST_DEVICE
0344   Value_t ValueAt(int level) const { return fPath[level]; }
0345
0346   // direct write access to the path
0347   // (no one should ever call this function unless you know what you are doing)
0348   // TODO: consider making this private + friend or so
0349   VECGEOM_FORCE_INLINE
0350   VECCORE_ATT_HOST_DEVICE
0351   void SetValueAt(int level, Value_t v) { fPath[level] = v; }
0352
0353   VECGEOM_FORCE_INLINE
0354   VECCORE_ATT_HOST_DEVICE
0355   void TopMatrix(Transformation3D &) const;
0356
0357 #ifdef VECGEOM_USE_CACHED_TRANSFORMATIONS
0358   /* @brief Update the cached top matrix */
0359   VECGEOM_FORCE_INLINE
0360   VECCORE_ATT_HOST_DEVICE
0361   void UpdateTopMatrix(Transformation3D *top_matrix = nullptr);
0362 #endif
0363
0364   // returning a "delta" transformation that can transform
0365   // coordinates given in reference frame of this->Top() to the reference frame of other->Top()
0366   // simply with otherlocalcoordinate = delta.Transform( thislocalcoordinate )
0367   VECCORE_ATT_HOST_DEVICE
0368   void DeltaTransformation(NavStatePath const &other, Transformation3D & /* delta */) const;
0369
0370   // VECGEOM_FORCE_INLINE
0371   VECCORE_ATT_HOST_DEVICE
0372   Vector3D<Precision> GlobalToLocal(Vector3D<Precision> const &) const;
0373
0374   VECCORE_ATT_HOST_DEVICE
0375   Vector3D<Precision> GlobalToLocal(Vector3D<Precision> const &, int tolevel) const;
0376
0377   VECCORE_ATT_HOST_DEVICE
0378   void TopMatrix(int tolevel, Transformation3D &) const;
0379
0380   VECGEOM_FORCE_INLINE
0381   VECCORE_ATT_HOST_DEVICE
0382   void Pop();
0383
0384   VECGEOM_FORCE_INLINE
0385   VECCORE_ATT_HOST_DEVICE
0386   int Distance(NavStatePath const &) const;
0387
0388   // returns a string representation of a (relative) sequence of operations/moves
0389   // that transforms this navigation state into the other navigation state
0390   // example:
0391   // state1 = /0/1/1/
0392   // state2 = /0/2/2/3
0393   // results in string
0394   // "/up/horiz/1/down/2/down/3" with 4 operations "up", "horiz", "down", "down"
0395   // the sequence of moves is the following
0396   // up: /0/1/1 --> /0/1/
0397   // horiz/1 : 0/1 --> /0/2 ( == /0/(1+1) )   "we are hopping from daughter 1 to 2 (which corresponds to a step of 1)"
0398   // down/2 : /0/2 --> /0/2/2   "going further down 2nd daughter"
0399   // down/3 : /0/2/2/3 --> /0/2/2/3  "going further down 2nd daughter"
0400   std::string RelativePath(NavStatePath const & /*other*/) const;
0401
0402   // clear all information
0403   VECGEOM_FORCE_INLINE
0404   VECCORE_ATT_HOST_DEVICE
0405   void Clear();
0406
0407   VECCORE_ATT_HOST_DEVICE
0408   void Print() const;
0409
0410   VECGEOM_FORCE_INLINE
0411   VECCORE_ATT_HOST_DEVICE
0412   void Dump() const;
0413
0414   VECGEOM_FORCE_INLINE
0415   VECCORE_ATT_HOST_DEVICE
0416   bool HasSamePathAsOther(NavStatePath const &other) const
0417   {
0418     if (other.fCurrentLevel != fCurrentLevel) return false;
0419     for (int i = fCurrentLevel - 1; i >= 0; --i) {
0420       if (fPath[i] != other.fPath[i]) return false;
0421     }
0422     return true;
0423   }
0424
0425   void printValueSequence(std::ostream & = std::cerr) const;
0426
0427   // calculates a checksum along the path
0428   // can be used (as a quick criterion) to see whether 2 states are same
0429   unsigned long getCheckSum() const
0430   {
0431     unsigned long s = 0;
0432     for (int i = 0; i < fCurrentLevel; ++i) {
0433       s += (unsigned long)(ValueAt(i) + 1); // + 1 as offset otherwise may not distinguish TOP level and OUTSIDE level
0434     }
0435     return s;
0436   }
0437
0438   /**
0439    * returns the number of FILLED LEVELS such that
0440    * state.GetNode( state.GetLevel() ) == state.Top()
0441    */
0442   VECGEOM_FORCE_INLINE
0443   VECCORE_ATT_HOST_DEVICE
0444   unsigned char GetLevel() const { return fCurrentLevel - 1; }
0445
0446   /**
0447     function returning whether the point (current navigation state) is outside the detector setup
0448   */
0449   VECGEOM_FORCE_INLINE
0450   VECCORE_ATT_HOST_DEVICE
0451   bool IsOutside() const { return !(fCurrentLevel > 0); }
0452
0453   VECGEOM_FORCE_INLINE
0454   VECCORE_ATT_HOST_DEVICE
0455   bool IsOnBoundary() const { return fOnBoundary; }
0456
0457 #ifdef VECGEOM_USE_CACHED_TRANSFORMATIONS
0458   VECGEOM_FORCE_INLINE
0459   VECCORE_ATT_HOST_DEVICE
0460   bool IsMatrixCached() const { return fCacheM; }
0461 #endif
0462
0463   VECGEOM_FORCE_INLINE
0464   VECCORE_ATT_HOST_DEVICE
0465   void SetBoundaryState(bool b) { fOnBoundary = b; }
0466
0467   VECGEOM_FORCE_INLINE
0468   VECCORE_ATT_HOST_DEVICE
0469   short GetCacheValue() const { return fCache; }
0470
0471   VECGEOM_FORCE_INLINE
0472   VECCORE_ATT_HOST_DEVICE
0473   void SetCacheValue(short v) { fCache = v; }
0474
0475 }; // end of class
0476
0477 VECCORE_ATT_HOST_DEVICE
0478 NavStatePath &NavStatePath::operator=(NavStatePath const &rhs)
0479 {
0480   if (this != &rhs) {
0481     fCurrentLevel = rhs.fCurrentLevel;
0482     fCache        = rhs.fCache;
0483     fOnBoundary   = rhs.fOnBoundary;
0484 #ifdef VECGEOM_USE_CACHED_TRANSFORMATIONS
0485     fCacheM   = rhs.fCacheM;
0486     fTopTrans = rhs.fTopTrans;
0487 #endif
0488     // Use memcpy.  Potential truncation if this is smaller than rhs.
0489     fPath = rhs.fPath;
0490   }
0491   return *this;
0492 }
0493
0494 // private implementation of standard constructor
0495 VECCORE_ATT_HOST_DEVICE
0496 NavStatePath::NavStatePath(size_t nvalues) : fCurrentLevel(0), fCache(-1), fOnBoundary(false), fPath(nvalues)
0497 {
0498   // clear the buffer
0499   std::memset(fPath.GetValues(), 0, nvalues * sizeof(NavStateIndex_t));
0500 }
0501
0502 VECCORE_ATT_HOST_DEVICE
0503 NavStatePath::~NavStatePath() {}
0504
0505 VECCORE_ATT_HOST_DEVICE
0506 void NavStatePath::Pop()
0507 {
0508   if (fCurrentLevel > 0) {
0509     fCurrentLevel--;
0510     // note that we are not invalidating the "popped volume" here
0511     // in order to be able to query the last "exited volume" later
0512     fCache = -1;
0513 #ifdef VECGEOM_USE_CACHED_TRANSFORMATIONS
0514     fCacheM = false;
0515 #endif
0516   }
0517 }
0518
0519 VECCORE_ATT_HOST_DEVICE
0520 void NavStatePath::Clear()
0521 {
0522   fCurrentLevel = 0;
0523   fOnBoundary   = false;
0524   fCache        = -1;
0525 #ifdef VECGEOM_USE_CACHED_TRANSFORMATIONS
0526   fCacheM = true;
0527 #endif
0528 }
0529
0530 VECCORE_ATT_HOST_DEVICE
0531 void NavStatePath::Push(VPlacedVolume const *v)
0532 {
0533 #ifdef DEBUG
0534   VECGEOM_ASSERT(fCurrentLevel < GetMaxLevel());
0535 #endif
0536   fPath[fCurrentLevel++] = ToIndex(v);
0537 #ifdef VECGEOM_USE_CACHED_TRANSFORMATIONS
0538   fCacheM = false;
0539 #endif
0540 }
0541
0542 // Allow pushing by child index
0543 VECCORE_ATT_HOST_DEVICE
0544 void NavStatePath::Push(unsigned short child)
0545 {
0546   if (fCurrentLevel > 0) {
0547     auto top = ToPlacedVolume(fPath[fCurrentLevel - 1]);
0548 #ifdef DEBUG
0549     VECGEOM_ASSERT(child < top->GetDaughters().size());
0550 #endif
0551     fPath[fCurrentLevel++] = ToIndex(top->GetDaughters().operator[](child));
0552 #ifdef VECGEOM_USE_CACHED_TRANSFORMATIONS
0553     fCacheM = false;
0554 #endif
0555   }
0556 }
0557
0558 VECCORE_ATT_HOST_DEVICE
0559 void NavStatePath::PushIndexType(NavStateIndex_t v)
0560 {
0561 #ifdef DEBUG
0562   VECGEOM_ASSERT(fCurrentLevel < GetMaxLevel());
0563 #endif
0564   fPath[fCurrentLevel++] = v;
0565 #ifdef VECGEOM_USE_CACHED_TRANSFORMATIONS
0566   fCacheM = false;
0567 #endif
0568 }
0569
0570 VECCORE_ATT_HOST_DEVICE
0571 VPlacedVolume const *NavStatePath::Top() const
0572 {
0573   return (fCurrentLevel > 0) ? ToPlacedVolume(fPath[fCurrentLevel - 1]) : nullptr;
0574 }
0575
0576 // calculates the global matrix to transform from global coordinates
0577 // to the frame of the top volume in the state
0578 // input: a reference to a transformation object ( which should be initialized to identity )
0579 VECGEOM_FORCE_INLINE
0580 VECCORE_ATT_HOST_DEVICE
0581 void NavStatePath::TopMatrix(Transformation3D &global_matrix) const
0582 {
0583   // this could be actually cached in case the path does not change ( particle stays inside a volume )
0584 #ifdef VECGEOM_USE_CACHED_TRANSFORMATIONS
0585   if (fCacheM) {
0586     global_matrix = fTopTrans;
0587     return;
0588   }
0589 #endif
0590   for (int i = fCurrentLevel - 1; i > 0; --i) {
0591     global_matrix *= *(ToPlacedVolume(fPath[i])->GetTransformation());
0592   }
0593 }
0594
0595 #ifdef VECGEOM_USE_CACHED_TRANSFORMATIONS
0596 // Update the cached top matrix
0597 VECGEOM_FORCE_INLINE
0598 VECCORE_ATT_HOST_DEVICE
0599 void NavStatePath::UpdateTopMatrix(Transformation3D *top_matrix)
0600 {
0601   // Update cached top transformation
0602   if (fCacheM) return;
0603   if (top_matrix) {
0604     fTopTrans = *top_matrix;
0605   } else {
0606     // need to recompute
0607     fTopTrans.Clear();
0608     TopMatrix(fTopTrans);
0609   }
0610   // Flag matrix as cached
0611   fCacheM = true;
0612 }
0613 #endif
0614
0615 VECGEOM_FORCE_INLINE
0616 VECCORE_ATT_HOST_DEVICE
0617 void NavStatePath::Dump() const
0618 {
0619   const unsigned int *ptr = (const unsigned int *)this;
0620 #pragma GCC diagnostic push
0621 #pragma GCC diagnostic ignored "-Wcast-qual"
0622   printf("NavState::Dump(): data: %p(%zu) : %p(%zu) : %p(%zu)\n", (void *)&fCurrentLevel, sizeof(fCurrentLevel),
0623          (void *)&fOnBoundary, sizeof(fOnBoundary), (void *)&fPath, sizeof(fPath));
0624   for (unsigned int i = 0; i < 20; ++i) {
0625     printf("%p: ", (void *)ptr);
0626     for (unsigned int j = 0; j < 8; ++j) {
0627       printf(" %08x ", *ptr);
0628       ptr++;
0629     }
0630     printf("\n");
0631   }
0632 #pragma GCC diagnostic pop
0633 }
0634
0635 /**
0636  * encodes the geometry path as a concatenated string of ( Value_t ) present in fPath
0637  */
0638 inline void NavStatePath::printValueSequence(std::ostream &stream) const
0639 {
0640   for (int i = 0; i < fCurrentLevel; ++i) {
0641     stream << "/" << fPath[i] << "(" << At(i)->GetLabel() << ")";
0642   }
0643 }
0644
0645 /**
0646  * calculates if other navigation state takes a different branch in geometry path or is on same branch
0647  * ( two states are on same branch if one can connect the states just by going upwards or downwards ( or do nothing ))
0648  */
0649 VECGEOM_FORCE_INLINE
0650 VECCORE_ATT_HOST_DEVICE
0651 int NavStatePath::Distance(NavStatePath const &other) const
0652 {
0653   int lastcommonlevel = -1;
0654   int maxlevel        = Min(GetCurrentLevel(), other.GetCurrentLevel());
0655
0656   //  algorithm: start on top and go down until paths split
0657   for (int i = 0; i < maxlevel; i++) {
0658     if (this->At(i) == other.At(i)) {
0659       lastcommonlevel = i;
0660     } else {
0661       break;
0662     }
0663   }
0664
0665   return (GetCurrentLevel() - lastcommonlevel) + (other.GetCurrentLevel() - lastcommonlevel) - 2;
0666 }
0667
0668 inline void NavStatePath::ConvertToGPUPointers()
0669 {
0670 #if !defined(VECCORE_CUDA) && defined(VECGEOM_ENABLE_CUDA) && !defined(VECGEOM_USE_INDEXEDNAVSTATES)
0671   for (int i = 0; i < fCurrentLevel; ++i) {
0672     auto *pvol = vecgeom::CudaManager::Instance().LookupPlaced(ToPlacedVolume(fPath[i])).GetPtr();
0673     fPath[i]   = ToIndex(pvol);
0674   }
0675 #endif
0676 }
0677
0678 inline void NavStatePath::ConvertToCPUPointers()
0679 {
0680 #if !defined(VECCORE_CUDA) && defined(VECGEOM_ENABLE_CUDA) && !defined(VECGEOM_USE_INDEXEDNAVSTATES)
0681   for (int i = 0; i < fCurrentLevel; ++i)
0682     fPath[i] = ToIndex(vecgeom::CudaManager::Instance().LookupPlacedCPUPtr((const void *)ToPlacedVolume(fPath[i])));
0683 #endif
0684 }
0685 } // namespace VECGEOM_IMPL_NAMESPACE
0686 } // namespace vecgeom
0687
0688 #if defined(GCC_DIAG_POP_NEEDED)
0689 #pragma GCC diagnostic pop
0690 #undef GCC_DIAG_POP_NEEDED
0691 #endif
0692
0693 #endif // VECGEOM_NAVIGATION_NAVSTATEPATH_H_