VecGeom/navigation/NavStatePool.h

0001 /*
0002  * NavStatePool.h
0003  *
0004  *  Created on: 14.11.2014
0005  *      Author: swenzel
0006  */
0007
0008 #ifndef NAVSTATEPOOL_H_
0009 #define NAVSTATEPOOL_H_
0010
0011 #include "VecGeom/base/Config.h"
0012 #include "VecGeom/base/Cuda.h"
0013 #include "VecGeom/base/Global.h"
0014 #include "VecGeom/base/Assert.h"
0015 #include "VecGeom/navigation/NavigationState.h"
0016 #include "VecGeom/management/Logger.h"
0017 #include <stdexcept>
0018
0019 #ifdef VECGEOM_ENABLE_CUDA
0020 #include "VecGeom/management/CudaManager.h"
0021 #endif
0022 #ifdef VECGEOM_CUDA_INTERFACE
0023 #include "VecGeom/backend/cuda/Interface.h"
0024 #endif
0025
0026 // a fixed (runtime) size  "array" or contiguous
0027 // memory pool of navigation states
0028 // testing some ideas to copy to gpu
0029 // it is supposed to be long-lived ( it has some initialization time overhead because it allocates the
0030 // GPU pointer at startup
0031
0032 #include <iostream>
0033 #include <fstream>
0034
0035 namespace vecgeom {
0036
0037 VECGEOM_DEVICE_FORWARD_DECLARE(template <typename Type> class SOA3D;);
0038 VECGEOM_HOST_FORWARD_DECLARE(class NavStatePool;);
0039 VECGEOM_DEVICE_FORWARD_DECLARE(class NavStatePool;);
0040 inline namespace VECGEOM_IMPL_NAMESPACE {
0041
0042 VECCORE_ATT_HOST_DEVICE
0043 VECCORE_FORCE_INLINE
0044 NavigationState const *GetNavigationState(int i, const char *buffer, int depth)
0045 {
0046   return reinterpret_cast<NavigationState const *>(buffer + NavigationState::SizeOfInstanceAlignAware(depth) * i);
0047 }
0048
0049 VECCORE_ATT_HOST_DEVICE
0050 VECCORE_FORCE_INLINE
0051 NavigationState *GetNavigationState(int i, char *buffer, int depth)
0052 {
0053   return reinterpret_cast<NavigationState *>(buffer + NavigationState::SizeOfInstanceAlignAware(depth) * i);
0054 }
0055
0056 class NavStatePoolView {
0057 public:
0058   VECCORE_ATT_HOST_DEVICE
0059   NavStatePoolView(char *buffer, int depth, int capacity) : fCapacity(capacity), fDepth(depth), fBuffer(buffer) {}
0060
0061   VECCORE_ATT_HOST_DEVICE
0062   NavigationState *operator[](int i)
0063   {
0064     VECGEOM_ASSERT(i < fCapacity);
0065     return GetNavigationState(i, fBuffer, fDepth);
0066   }
0067
0068   VECCORE_ATT_HOST_DEVICE
0069   NavigationState const *operator[](int i) const
0070   {
0071     VECGEOM_ASSERT(i < fCapacity);
0072     return GetNavigationState(i, fBuffer, fDepth);
0073   }
0074
0075   VECCORE_ATT_HOST_DEVICE
0076   int Capacity() const { return fCapacity; }
0077
0078   VECCORE_ATT_HOST_DEVICE
0079   int Depth() const { return fDepth; }
0080
0081   VECCORE_ATT_HOST_DEVICE
0082   int IsValid() const { return fBuffer && fCapacity > 0 && fDepth > 0; }
0083
0084 private:         // members
0085   int fCapacity; // Allocated size of the container.
0086   int fDepth;    // depth of the navigation objects to cover
0087   char *fBuffer; // the memory buffer in which we place states
0088 }; // end class
0089
0090 class NavStatePool {
0091
0092 public:
0093   NavStatePool(int size, int depth)
0094       : fCapacity(size), fDepth(depth), fBuffer(new char[NavigationState::SizeOfInstanceAlignAware(depth) * size]),
0095         fGPUPointer(NULL)
0096   {
0097
0098 #ifdef VECGEOM_CUDA_INTERFACE
0099     VECGEOM_DEVICE_API_CALL(Malloc(&fGPUPointer, NavigationState::SizeOfInstanceAlignAware(depth) * size));
0100 #endif
0101     // now create the states
0102     for (int i = 0; i < (int)fCapacity; ++i) {
0103       NavigationState::MakeInstanceAt(depth, fBuffer + NavigationState::SizeOfInstanceAlignAware(depth) * i);
0104     }
0105   }
0106
0107   ~NavStatePool()
0108   {
0109 #ifdef VECGEOM_CUDA_INTERFACE
0110     try {
0111       VECGEOM_DEVICE_API_CALL(Free(fGPUPointer));
0112     } catch (std::runtime_error const &e) {
0113       VECGEOM_LOG(error) << e.what();
0114     }
0115 #endif
0116     delete[] fBuffer;
0117   }
0118 #if !defined(VECCORE_CUDA) && defined(VECGEOM_ENABLE_CUDA)
0119   void CopyToGpu();
0120   void CopyFromGpu();
0121 #endif
0122
0123   // quick and dirty serialization and deserialization
0124   void ToFile(std::string filename) const
0125   {
0126 #ifdef VECGEOM_USE_INDEXEDNAVSTATES
0127     std::ofstream outfile(filename, std::ios::binary);
0128     outfile.write(reinterpret_cast<const char *>(&fCapacity), sizeof(fCapacity));
0129     outfile.write(reinterpret_cast<const char *>(&fDepth), sizeof(fDepth));
0130     outfile.write(reinterpret_cast<char *>(fBuffer), fCapacity * NavigationState::SizeOfInstanceAlignAware(fDepth));
0131 #else
0132     std::cerr << "serializing pointer based navstates not supported \n";
0133 #endif
0134   }
0135
0136   static void ReadDepthAndCapacityFromFile(std::string filename, int &cap, int &dep)
0137   {
0138     std::ifstream fin(filename, std::ios::binary);
0139     fin.read(reinterpret_cast<char *>(&cap), sizeof(cap));
0140     fin.read(reinterpret_cast<char *>(&dep), sizeof(dep));
0141   }
0142
0143   // return number of elements read or -1 if failure
0144   int FromFile(std::string filename)
0145   {
0146 #ifdef VECGEOM_USE_INDEXEDNAVSTATES
0147     // assumes existing NavStatePool object
0148     decltype(fCapacity) cap;
0149     decltype(fDepth) dep;
0150     std::ifstream fin(filename, std::ios::binary);
0151     if (!fin) return -1;
0152     fin.read(reinterpret_cast<char *>(&cap), sizeof(cap));
0153     if (!fin) return -2;
0154     fin.read(reinterpret_cast<char *>(&dep), sizeof(dep));
0155     if (!fin) return -2;
0156     if (cap != fCapacity || dep != fDepth) std::cerr << " warning: reading from navstate with different size\n";
0157     fin.read(reinterpret_cast<char *>(fBuffer), fCapacity * NavigationState::SizeOfInstanceAlignAware(fDepth));
0158     if (!fin) return -3;
0159 #else
0160     std::cerr << "serializing pointer based navstates not supported \n";
0161 #endif
0162     return fCapacity;
0163   }
0164
0165   VECCORE_ATT_HOST_DEVICE
0166   NavigationState *operator[](int i) { return GetNavigationState(i, fBuffer, fDepth); }
0167
0168   VECCORE_ATT_HOST_DEVICE
0169   NavigationState const *operator[](int i) const { return GetNavigationState(i, fBuffer, fDepth); }
0170
0171   // convert/init this to a plain NavigationState** array
0172   // so that array[0] points to the first state in the NavStatePool, etc
0173   // this method also allocates memory; array should be a nullptr initially
0174   // this is a convenience function
0175   VECCORE_ATT_HOST_DEVICE
0176   void ToPlainPointerArray(NavigationState const **&array) const
0177   {
0178     array = new NavigationState const *[fCapacity];
0179     for (int i = 0; i < fCapacity; ++i) {
0180       array[i] = (*this)[i];
0181     }
0182   }
0183
0184   // dito for the non-const version
0185   VECCORE_ATT_HOST_DEVICE
0186   void ToPlainPointerArray(NavigationState **&array)
0187   {
0188     array = new NavigationState *[fCapacity];
0189     for (int i = 0; i < fCapacity; ++i) {
0190       array[i] = (*this)[i];
0191     }
0192   }
0193
0194   void Print() const
0195   {
0196     for (int i = 0; i < fCapacity; ++i)
0197       (*this)[i]->Print();
0198   }
0199
0200   void *GetGPUPointer() const { return fGPUPointer; }
0201
0202   int capacity() const { return fCapacity; }
0203
0204 private: // protected methods
0205 #ifdef VECGEOM_ENABLE_CUDA
0206          // This constructor used to build NavStatePool at the GPU.  BufferGPU
0207   VECCORE_ATT_DEVICE
0208   NavStatePool(int size, int depth, char *fBufferGPU)
0209       : fCapacity(size), fDepth(depth), fBuffer(fBufferGPU), fGPUPointer(NULL)
0210   {
0211   }
0212 #endif
0213
0214 private:         // members
0215   int fCapacity; // the number of states in the pool
0216   int fDepth;    // depth of the navigation objects to cover
0217   char *fBuffer; // the memory buffer in which we place states
0218
0219   // assume it keeps a GPU pointer directly
0220   // the target of the copy operation
0221   void *fGPUPointer;
0222
0223 }; // end class
0224
0225 // an implementation of the CopyOperation could be as follows
0226 #if !defined(VECCORE_CUDA) && defined(VECGEOM_ENABLE_CUDA)
0227 inline void NavStatePool::CopyToGpu()
0228 {
0229
0230   // modify content temporarily to convert CPU pointers to GPU pointers
0231   NavigationState *state;
0232   for (int i = 0; i < fCapacity; ++i) {
0233     state = operator[](i);
0234     state->ConvertToGPUPointers();
0235   }
0236
0237   // we also have to fix the fPath pointers
0238
0239   // copy
0240   vecgeom::CopyToGpu((void *)fBuffer, fGPUPointer, fCapacity * NavigationState::SizeOfInstanceAlignAware(fDepth));
0241   // CudaAssertError( cudaMemcpy(fGPUPointer, (void*)fBuffer, fCapacity*NavigationState::SizeOf(fDepth),
0242   // VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)) );
0243
0244   // modify back pointers
0245   for (int i = 0; i < fCapacity; ++i) {
0246     state = operator[](i);
0247     state->ConvertToCPUPointers();
0248   }
0249
0250   // now some kernel can be launched on GPU side
0251 } // end CopyFunction
0252
0253 inline void NavStatePool::CopyFromGpu()
0254 {
0255   // this does not work
0256   // modify content temporarily to convert CPU pointers to GPU pointers
0257
0258   // std::cerr << "Starting to COPY" << std::endl;
0259   // std::cerr << "GPU pointer " << fGPUPointer << std::endl;
0260   vecgeom::CopyFromGpu(fGPUPointer, (void *)fBuffer, fCapacity * NavigationState::SizeOfInstanceAlignAware(fDepth));
0261
0262   NavigationState *state;
0263   for (int i = 0; i < fCapacity; ++i) {
0264     state = operator[](i);
0265     state->ConvertToCPUPointers();
0266   }
0267 } // end CopyFunction
0268 #endif
0269 } // namespace VECGEOM_IMPL_NAMESPACE
0270 } // namespace vecgeom
0271
0272 #endif /* NAVSTATEPOOL_H_ */