surfaces/cuda/BrepCudaManager.h

0001 #ifndef VECGEOM_SURFACE_BREPCUDAMANAGER_H_
0002 #define VECGEOM_SURFACE_BREPCUDAMANAGER_H_
0003
0004 #include <VecGeom/surfaces/SurfData.h>
0005 #include "VecGeom/base/AABB.h"
0006 #include "VecGeom/base/BVH.h"
0007 #include "VecGeom/volumes/VolumeTree.h"
0008 #include "VecGeom/management/Logger.h"
0009 #include "VecGeom/base/Assert.h"
0010
0011 namespace vgbrep {
0012
0013 template <typename Real_t>
0014 __global__ void FinishBVHCopy(vecgeom::BVH<Real_t> *dBVH, int *dPrimId, int *dOffset, int *dNChild,
0015                               vecgeom::AABB<typename SurfData<Real_t>::Real_b> *dAABBs,
0016                               vecgeom::AABB<typename SurfData<Real_t>::Real_b> *dNodes)
0017 {
0018   if (dBVH == nullptr) {
0019     printf("Error: Null pointer 'dBVH' encountered in FinishBVHCopy\n");
0020     return;
0021   }
0022   if (dPrimId == nullptr) {
0023     printf("Error: Null pointer 'dPrimId' encountered in FinishBVHCopy\n");
0024     return;
0025   }
0026   if (dOffset == nullptr) {
0027     printf("Error: Null pointer 'dOffset' encountered in FinishBVHCopy\n");
0028     return;
0029   }
0030   if (dNChild == nullptr) {
0031     printf("Error: Null pointer 'dNChild' encountered in FinishBVHCopy\n");
0032     return;
0033   }
0034   if (dAABBs == nullptr) {
0035     printf("Error: Null pointer 'dAABBs' encountered in FinishBVHCopy\n");
0036     return;
0037   }
0038   if (dNodes == nullptr) {
0039     printf("Error: Null pointer 'dNodes' encountered in FinishBVHCopy\n");
0040     return;
0041   }
0042   dBVH->SetPointers(dPrimId, dOffset, dNChild, dAABBs, dNodes);
0043 }
0044
0045 template <typename Real_t>
0046 void CopyBVH(const vecgeom::BVH<Real_t> &hBVH, vecgeom::BVH<Real_t> *dBVH)
0047 {
0048   int *dPrimId;
0049   int *dOffset;
0050   int *dNChild;
0051   vecgeom::AABB<Real_t> *dNodes;
0052   vecgeom::AABB<Real_t> *dAABBs;
0053
0054   int rootNChild = hBVH.GetRootNChild();
0055   if (rootNChild <= 0) {
0056     std::ostringstream oss;
0057     oss << "Invalid number of root children: " << rootNChild;
0058     throw std::logic_error(oss.str());
0059   }
0060   int nodes = (2 << hBVH.GetDepth()) - 1;
0061   if (nodes <= 0) {
0062     std::ostringstream oss;
0063     oss << "Invalid number of nodes: " << nodes;
0064     throw std::logic_error(oss.str());
0065   }
0066
0067   VECGEOM_DEVICE_API_CALL(Malloc(&dPrimId, hBVH.GetRootNChild() * sizeof(int)));
0068   VECGEOM_DEVICE_API_CALL(Malloc(&dOffset, nodes * sizeof(int)));
0069   VECGEOM_DEVICE_API_CALL(Malloc(&dNChild, nodes * sizeof(int)));
0070   VECGEOM_DEVICE_API_CALL(Malloc(&dNodes, nodes * sizeof(vecgeom::AABB<Real_t>)));
0071   VECGEOM_DEVICE_API_CALL(Malloc(&dAABBs, hBVH.GetRootNChild() * sizeof(vecgeom::AABB<Real_t>)));
0072
0073   // Ensure pointers are not null after allocation
0074   if (!dPrimId || !dAABBs || !dOffset || !dNChild || !dNodes) {
0075     throw std::runtime_error("Memory allocation failed: One or more pointers are null.");
0076   }
0077
0078   VECGEOM_DEVICE_API_CALL(Memcpy(dPrimId, hBVH.GetPrimId(), hBVH.GetRootNChild() * sizeof(int),
0079                                  VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0080   VECGEOM_DEVICE_API_CALL(
0081       Memcpy(dOffset, hBVH.GetOffset(), nodes * sizeof(int), VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0082   VECGEOM_DEVICE_API_CALL(
0083       Memcpy(dNChild, hBVH.GetNChild(), nodes * sizeof(int), VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0084   VECGEOM_DEVICE_API_CALL(Memcpy(dNodes, hBVH.GetNodes(), nodes * sizeof(vecgeom::AABB<Real_t>),
0085                                  VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0086   VECGEOM_DEVICE_API_CALL(Memcpy(dAABBs, hBVH.GetAABBs(), hBVH.GetRootNChild() * sizeof(vecgeom::AABB<Real_t>),
0087                                  VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0088
0089   // Adjust pointers in the GPU instance
0090   FinishBVHCopy<<<1, 1>>>(dBVH, dPrimId, dOffset, dNChild, dAABBs, dNodes);
0091 }
0092
0093 // This function sets the correct pointers on device memory in the data structures that were copied
0094 static __global__ void FinishVolumeTreeTransfer(vecgeom::VolumeTree *volumeTree, long shift_children)
0095 {
0096   volumeTree->Relocate(shift_children);
0097   vecgeom::globaldevicegeomdata::gVolumeTree = volumeTree;
0098 }
0099
0100 // This function sets the correct pointers on device memory in the data structures that were copied
0101 template <typename Real_t>
0102 static __global__ void BrepCudaManagerFinishTransfer(SurfData<Real_t> *surfData)
0103 {
0104   int *current, *current_exiting_surface, *current_entering_surface, *current_entering_surface_pvol,
0105       *current_entering_surface_pvol_trans, *current_entering_surface_lvol_id, *current_daughter_pvol_id,
0106       *current_daughter_pvol_trans;
0107   logic_int *current_logic;
0108   globaldevicesurfdata::gSurfDataDevice<Real_t> = surfData;
0109
0110   // Write pointers into fShells[i].fSurfaces, fShells[i].fLogic, fShells[i].fShellEnteringSurfaceList,
0111   // fShells[i].fShellEnteringSurfaceTransList, fShells[i].fShellEnteringSurfacePvolList
0112   current                             = surfData->fSurfShellList;
0113   current_logic                       = surfData->fLogicList;
0114   current_exiting_surface             = surfData->fShellExitingSurfaceList;
0115   current_entering_surface            = surfData->fShellEnteringSurfaceList;
0116   current_entering_surface_pvol       = surfData->fShellEnteringSurfacePvolList;
0117   current_entering_surface_pvol_trans = surfData->fShellEnteringSurfacePvolTransList;
0118   current_entering_surface_lvol_id    = surfData->fShellEnteringSurfaceLvolIdList;
0119   current_daughter_pvol_id            = surfData->fShellDaughterPvolIdList;
0120   current_daughter_pvol_trans         = surfData->fShellDaughterPvolTransList;
0121   for (int i = 0; i < surfData->fNshells; i++) {
0122     surfData->fShells[i].fSurfaces = current;
0123     current += surfData->fShells[i].fNsurf;
0124     surfData->fShells[i].fLogic.data_ = current_logic;
0125     current_logic += surfData->fShells[i].fLogic.size();
0126     surfData->fShells[i].fExitingSurfaces = current_exiting_surface;
0127     current_exiting_surface += surfData->fShells[i].fNExitingSurfaces;
0128     surfData->fShells[i].fEnteringSurfaces = current_entering_surface;
0129     current_entering_surface += surfData->fShells[i].fNEnteringSurfaces;
0130     surfData->fShells[i].fEnteringSurfacesPvol = current_entering_surface_pvol;
0131     current_entering_surface_pvol += surfData->fShells[i].fNEnteringSurfaces;
0132     surfData->fShells[i].fEnteringSurfacesPvolTrans = current_entering_surface_pvol_trans;
0133     current_entering_surface_pvol_trans += surfData->fShells[i].fNEnteringSurfaces;
0134     surfData->fShells[i].fEnteringSurfacesLvolIds = current_entering_surface_lvol_id;
0135     current_entering_surface_lvol_id += surfData->fShells[i].fNEnteringSurfaces;
0136     surfData->fShells[i].fDaughterPvolIds = current_daughter_pvol_id;
0137     current_daughter_pvol_id += surfData->fShells[i].fNDaughterPvols;
0138     surfData->fShells[i].fDaughterPvolTrans = current_daughter_pvol_trans;
0139     current_daughter_pvol_trans += surfData->fShells[i].fNDaughterPvols;
0140   }
0141
0142   // Write pointers into fCommonSurfaces[i].f{Left,Right}Side.fSurfaces
0143   current = surfData->fSides;
0144   for (int i = 0; i < surfData->fNcommonSurf; i++) {
0145     surfData->fCommonSurfaces[i].fLeftSide.fSurfaces = current;
0146     current += surfData->fCommonSurfaces[i].fLeftSide.fNsurf;
0147     surfData->fCommonSurfaces[i].fRightSide.fSurfaces = current;
0148     current += surfData->fCommonSurfaces[i].fRightSide.fNsurf;
0149   }
0150
0151   // Write pointers into fSideDivisions[i].fSlices and fSlices[i].fCandidates
0152   auto currentSlice = surfData->fSlices;
0153   current           = surfData->fSliceCandidates;
0154   for (int i = 0; i < surfData->fNsideDivisions; i++) {
0155     surfData->fSideDivisions[i].fSlices = currentSlice;
0156     for (int j = 0; j < surfData->fSideDivisions[i].fNslices; ++j) {
0157       currentSlice->fCandidates = current;
0158       current += currentSlice->fNcand;
0159       currentSlice++;
0160     }
0161   }
0162
0163   // Write pointers into fCandidates[i].{fCandidates,fFrameInd, fSides}
0164   current = surfData->fCandList;
0165   for (int i = 0; i < surfData->fNStates; i++) {
0166     surfData->fCandidates[i].fCandidates = current;
0167     // Move the pointer to the start of the Frame index list
0168     int ncand = surfData->fCandidates[i].fNcand;
0169     current += ncand;
0170     surfData->fCandidates[i].fFrameInd = current;
0171     // Move the pointer to the start of the sides index list
0172     current += ncand;
0173     surfData->fCandidates[i].fSides = reinterpret_cast<char *>(current);
0174     // Move the pointer to the start of the next Candidate index list
0175     int add_one = ((ncand * sizeof(char)) % sizeof(int)) > 0 ? 1 : 0;
0176     current += ncand * sizeof(char) / sizeof(int) + add_one;
0177   }
0178 }
0179
0180 // Manager class for synchronizing the surface data to the GPU.
0181 template <typename Real_t>
0182 class BrepCudaManager {
0183   using SurfData_t = SurfData<Real_t>;
0184
0185   SurfData_t fSurfDataStaging;               ///< Host memory to stage data for the GPU
0186   SurfData_t *fSurfData{nullptr};            ///< Device pointer to the data structure
0187   vecgeom::VolumeTree fVolumeTreeStaging;    ///< Host memory to stage the volume tree for the GPU
0188   vecgeom::VolumeTree *fVolumeTree{nullptr}; ///< Device pointer to the volume tree
0189
0190 public:
0191   static BrepCudaManager &Instance()
0192   {
0193     static BrepCudaManager instance;
0194     return instance;
0195   }
0196
0197   const SurfData_t *GetDevicePtr() const { return fSurfData; }
0198   const vecgeom::VolumeTree *GetVolumeTreeDevicePtr() const { return fVolumeTree; }
0199
0200   void TransferVolumeTree(const vecgeom::VolumeTree &volumeTree)
0201   {
0202     fVolumeTreeStaging  = volumeTree;
0203     size_t sizeLogical  = volumeTree.fNlogical * sizeof(vecgeom::LogicalId);
0204     size_t sizePlaced   = volumeTree.fNplaced * sizeof(vecgeom::PlacedId);
0205     size_t sizeChildren = volumeTree.fNplacedC * sizeof(vecgeom::PlacedId);
0206     // Allocate the buffer on device
0207     VECGEOM_DEVICE_API_CALL(Malloc(&fVolumeTreeStaging.fLogical, sizeLogical));
0208     VECGEOM_DEVICE_API_CALL(Malloc(&fVolumeTreeStaging.fPlaced, sizePlaced));
0209     VECGEOM_DEVICE_API_CALL(Malloc(&fVolumeTreeStaging.fChildren, sizeChildren));
0210     VECGEOM_DEVICE_API_CALL(Memcpy(fVolumeTreeStaging.fLogical, volumeTree.fLogical, sizeLogical,
0211                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0212     VECGEOM_DEVICE_API_CALL(Memcpy(fVolumeTreeStaging.fPlaced, volumeTree.fPlaced, sizePlaced,
0213                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0214     VECGEOM_DEVICE_API_CALL(Memcpy(fVolumeTreeStaging.fChildren, volumeTree.fChildren, sizeChildren,
0215                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0216     // long shift_placed = reinterpret_cast<long>(fVolumeTreeStaging.fPlaced) -
0217     // reinterpret_cast<long>(volumeTree.fPlaced);
0218     long shift_children =
0219         reinterpret_cast<long>(fVolumeTreeStaging.fChildren) - reinterpret_cast<long>(volumeTree.fChildren);
0220
0221     // Now copy the staged data to the GPU
0222     VECGEOM_DEVICE_API_CALL(Malloc(&fVolumeTree, sizeof(vecgeom::VolumeTree)));
0223     VECGEOM_DEVICE_API_CALL(Memcpy(fVolumeTree, &fVolumeTreeStaging, sizeof(vecgeom::VolumeTree),
0224                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0225
0226     // Finally finish the transfer by calling a kernel to write some pointers
0227     FinishVolumeTreeTransfer<<<1, 1>>>(fVolumeTree, shift_children);
0228     VECGEOM_DEVICE_API_CALL(DeviceSynchronize());
0229     // The arrays in the staging area are device pointers, null them to avoid deletion
0230     fVolumeTreeStaging.fLogical  = nullptr;
0231     fVolumeTreeStaging.fPlaced   = nullptr;
0232     fVolumeTreeStaging.fChildren = nullptr;
0233   }
0234
0235   void TransferSurfData(const SurfData_t &surfData)
0236   {
0237     using Real_b = typename SurfData_t::Real_b;
0238     size_t sizeInBytes;
0239
0240     // Transfer the volume tree first
0241     auto const &volumeTree = vecgeom::VolumeTree::Instance();
0242     if (!volumeTree.fValid) {
0243       VECGEOM_LOG(critical) << "Volume tree is invalid";
0244       return;
0245     }
0246     TransferVolumeTree(volumeTree);
0247
0248     // Allocate and copy transformations
0249     fSurfDataStaging.fNvolTrans = surfData.fNvolTrans;
0250     sizeInBytes                 = sizeof(surfData.fPVolTrans[0]) * surfData.fNvolTrans;
0251     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fPVolTrans, sizeInBytes));
0252     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fPVolTrans, surfData.fPVolTrans, sizeInBytes,
0253                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0254
0255     // Allocate and copy surface data
0256     fSurfDataStaging.fNellip = surfData.fNellip;
0257     sizeInBytes              = sizeof(surfData.fEllipData[0]) * surfData.fNellip;
0258     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fEllipData, sizeInBytes));
0259     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fEllipData, surfData.fEllipData, sizeInBytes,
0260                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0261
0262     fSurfDataStaging.fNtorus = surfData.fNtorus;
0263     sizeInBytes              = sizeof(surfData.fTorusData[0]) * surfData.fNtorus;
0264     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fTorusData, sizeInBytes));
0265     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fTorusData, surfData.fTorusData, sizeInBytes,
0266                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0267
0268     fSurfDataStaging.fNarb4 = surfData.fNarb4;
0269     sizeInBytes             = sizeof(surfData.fArb4Data[0]) * surfData.fNarb4;
0270     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fArb4Data, sizeInBytes));
0271     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fArb4Data, surfData.fArb4Data, sizeInBytes,
0272                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0273
0274     // Allocate and copy volume shells
0275     fSurfDataStaging.fNshells = surfData.fNshells;
0276     sizeInBytes               = sizeof(surfData.fShells[0]) * surfData.fNshells;
0277     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fShells, sizeInBytes));
0278     VECGEOM_DEVICE_API_CALL(
0279         Memcpy(fSurfDataStaging.fShells, surfData.fShells, sizeInBytes, VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0280
0281     // Nota bene: the fShells[i].fSurfaces are backed by the following array
0282     // and set via BrepCudaManagerFinishTransfer.
0283     sizeInBytes = sizeof(surfData.fSurfShellList[0]) * surfData.fNlocalSurf;
0284     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fSurfShellList, sizeInBytes));
0285     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fSurfShellList, surfData.fSurfShellList, sizeInBytes,
0286                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0287
0288     // Nota bene: fShells[i].fLogic are backed by the following array
0289     // and set via BrepCudaManagerFinishTransfer.
0290     sizeInBytes = sizeof(logic_int) * surfData.fNlogic;
0291     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fLogicList, sizeInBytes));
0292     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fLogicList, surfData.fLogicList, sizeInBytes,
0293                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0294
0295     // Nota bene: fShells[i].fShellExiting/EnteringSurfaceList are backed by the following array
0296     // and set via BrepCudaManagerFinishTransfer.
0297     fSurfDataStaging.fNExitingSurfaces  = surfData.fNExitingSurfaces;
0298     fSurfDataStaging.fNEnteringSurfaces = surfData.fNEnteringSurfaces;
0299
0300     sizeInBytes = sizeof(surfData.fShellExitingSurfaceList[0]) * surfData.fNExitingSurfaces;
0301     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fShellExitingSurfaceList, sizeInBytes));
0302     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fShellExitingSurfaceList, surfData.fShellExitingSurfaceList,
0303                                    sizeInBytes, VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0304     sizeInBytes = sizeof(surfData.fShellEnteringSurfaceList[0]) * surfData.fNEnteringSurfaces;
0305     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fShellEnteringSurfaceList, sizeInBytes));
0306     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fShellEnteringSurfaceList, surfData.fShellEnteringSurfaceList,
0307                                    sizeInBytes, VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0308
0309     // Nota bene: fShells[i].fShellEnteringSurfacePvolList are backed by the following array
0310     // and set via BrepCudaManagerFinishTransfer.
0311     sizeInBytes = sizeof(surfData.fShellEnteringSurfacePvolList[0]) * surfData.fNEnteringSurfaces;
0312     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fShellEnteringSurfacePvolList, sizeInBytes));
0313     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fShellEnteringSurfacePvolList,
0314                                    surfData.fShellEnteringSurfacePvolList, sizeInBytes,
0315                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0316
0317     sizeInBytes = sizeof(surfData.fShellEnteringSurfacePvolTransList[0]) * surfData.fNEnteringSurfaces;
0318     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fShellEnteringSurfacePvolTransList, sizeInBytes));
0319     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fShellEnteringSurfacePvolTransList,
0320                                    surfData.fShellEnteringSurfacePvolTransList, sizeInBytes,
0321                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0322
0323     sizeInBytes = sizeof(surfData.fShellEnteringSurfaceLvolIdList[0]) * surfData.fNEnteringSurfaces;
0324     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fShellEnteringSurfaceLvolIdList, sizeInBytes));
0325     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fShellEnteringSurfaceLvolIdList,
0326                                    surfData.fShellEnteringSurfaceLvolIdList, sizeInBytes,
0327                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0328
0329     sizeInBytes = sizeof(surfData.fShellDaughterPvolIdList[0]) * surfData.fNPlacedVolumes;
0330     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fShellDaughterPvolIdList, sizeInBytes));
0331     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fShellDaughterPvolIdList, surfData.fShellDaughterPvolIdList,
0332                                    sizeInBytes, VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0333
0334     sizeInBytes = sizeof(surfData.fShellDaughterPvolTransList[0]) * surfData.fNPlacedVolumes;
0335     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fShellDaughterPvolTransList, sizeInBytes));
0336     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fShellDaughterPvolTransList, surfData.fShellDaughterPvolTransList,
0337                                    sizeInBytes, VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0338
0339     // Allocate space for the BVHs
0340     // Surface BVHs
0341     sizeInBytes = sizeof(surfData.fBVH[0]) * surfData.fNshells;
0342     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fBVH, sizeInBytes));
0343     VECGEOM_DEVICE_API_CALL(
0344         Memcpy(fSurfDataStaging.fBVH, surfData.fBVH, sizeInBytes, VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0345     // Solid BVHs
0346     sizeInBytes = sizeof(surfData.fBVHSolids[0]) * surfData.fNshells;
0347     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fBVHSolids, sizeInBytes));
0348     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fBVHSolids, surfData.fBVHSolids, sizeInBytes,
0349                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0350
0351     // Allocate and copy BVH members
0352     for (int i = 0; i < surfData.fNshells; ++i) {
0353       // Surface BVHs
0354       if (surfData.fShells[i].fNsurf > 0) { // this checks that the BVH is actually populated and not null
0355         auto const &hBVH = surfData.fBVH[i];
0356         auto dBVH        = &(fSurfDataStaging.fBVH[i]);
0357         CopyBVH<Real_b>(hBVH, dBVH);
0358       }
0359       // Solid BVHs
0360       if (surfData.fShells[i].fNEnteringSurfaces >
0361           0) { // If there are no entering surfaces, the volume has no daughters
0362         auto const &hBVH = surfData.fBVHSolids[i];
0363         auto dBVH        = &(fSurfDataStaging.fBVHSolids[i]);
0364         CopyBVH<Real_b>(hBVH, dBVH);
0365       }
0366     }
0367
0368     // Allocate and copy surfaces
0369     fSurfDataStaging.fNlocalSurf = surfData.fNlocalSurf;
0370     sizeInBytes                  = sizeof(surfData.fLocalSurf[0]) * surfData.fNlocalSurf;
0371     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fLocalSurf, sizeInBytes));
0372     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fLocalSurf, surfData.fLocalSurf, sizeInBytes,
0373                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0374
0375     fSurfDataStaging.fNglobalSurf = surfData.fNglobalSurf;
0376     sizeInBytes                   = sizeof(surfData.fFramedSurf[0]) * surfData.fNglobalSurf;
0377     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fFramedSurf, sizeInBytes));
0378     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fFramedSurf, surfData.fFramedSurf, sizeInBytes,
0379                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0380
0381     // Allocate and copy masks
0382     fSurfDataStaging.fNwindows = surfData.fNwindows;
0383     sizeInBytes                = sizeof(surfData.fWindowMasks[0]) * surfData.fNwindows;
0384     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fWindowMasks, sizeInBytes));
0385     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fWindowMasks, surfData.fWindowMasks, sizeInBytes,
0386                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0387
0388     fSurfDataStaging.fNrings = surfData.fNrings;
0389     sizeInBytes              = sizeof(surfData.fRingMasks[0]) * surfData.fNrings;
0390     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fRingMasks, sizeInBytes));
0391     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fRingMasks, surfData.fRingMasks, sizeInBytes,
0392                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0393
0394     fSurfDataStaging.fNzphis = surfData.fNzphis;
0395     sizeInBytes              = sizeof(surfData.fZPhiMasks[0]) * surfData.fNzphis;
0396     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fZPhiMasks, sizeInBytes));
0397     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fZPhiMasks, surfData.fZPhiMasks, sizeInBytes,
0398                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0399
0400     fSurfDataStaging.fNquads = surfData.fNquads;
0401     sizeInBytes              = sizeof(surfData.fQuadMasks[0]) * surfData.fNquads;
0402     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fQuadMasks, sizeInBytes));
0403     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fQuadMasks, surfData.fQuadMasks, sizeInBytes,
0404                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0405
0406     fSurfDataStaging.fNtriangs = surfData.fNtriangs;
0407     sizeInBytes                = sizeof(surfData.fTriangleMasks[0]) * surfData.fNtriangs;
0408     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fTriangleMasks, sizeInBytes));
0409     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fTriangleMasks, surfData.fTriangleMasks, sizeInBytes,
0410                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0411
0412     // Allocate and copy scene indices
0413     fSurfDataStaging.fNscenes = surfData.fNscenes;
0414     sizeInBytes               = sizeof(int) * surfData.fNscenes;
0415     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fSceneStartIndex, sizeInBytes));
0416     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fSceneTouchables, sizeInBytes));
0417     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fSceneStartIndex, surfData.fSceneStartIndex, sizeInBytes,
0418                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0419     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fSceneTouchables, surfData.fSceneTouchables, sizeInBytes,
0420                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0421
0422     // Allocate and copy common surfaces
0423     fSurfDataStaging.fNcommonSurf = surfData.fNcommonSurf;
0424     sizeInBytes                   = sizeof(surfData.fCommonSurfaces[0]) * surfData.fNcommonSurf;
0425     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fCommonSurfaces, sizeInBytes));
0426     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fCommonSurfaces, surfData.fCommonSurfaces, sizeInBytes,
0427                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0428
0429     // Nota bene: the fCommonSurfaces[i].f{Left,Right}Side.fSurfaces are backed
0430     // by the following array and set via BrepCudaManagerFinishTransfer.
0431     fSurfDataStaging.fNsides = surfData.fNsides;
0432     sizeInBytes              = sizeof(surfData.fSides[0]) * surfData.fNsides;
0433     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fSides, sizeInBytes));
0434     VECGEOM_DEVICE_API_CALL(
0435         Memcpy(fSurfDataStaging.fSides, surfData.fSides, sizeInBytes, VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0436
0437     // Allocate and copy side divisions
0438     // Nota bene: the fSideDivisions[i].fSlices and fSlices[i].fCandidates are backed
0439     // by the following arrays and set via BrepCudaManagerFinishTransfer.
0440     fSurfDataStaging.fNsideDivisions   = surfData.fNsideDivisions;
0441     fSurfDataStaging.fNslices          = surfData.fNslices;
0442     fSurfDataStaging.fNsliceCandidates = surfData.fNsliceCandidates;
0443     sizeInBytes                        = sizeof(surfData.fSideDivisions[0]) * surfData.fNsideDivisions;
0444     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fSideDivisions, sizeInBytes));
0445     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fSideDivisions, surfData.fSideDivisions, sizeInBytes,
0446                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0447     sizeInBytes = sizeof(surfData.fSlices[0]) * surfData.fNslices;
0448     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fSlices, sizeInBytes));
0449     VECGEOM_DEVICE_API_CALL(
0450         Memcpy(fSurfDataStaging.fSlices, surfData.fSlices, sizeInBytes, VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0451     sizeInBytes = sizeof(int) * surfData.fNsliceCandidates;
0452     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fSliceCandidates, sizeInBytes));
0453     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fSliceCandidates, surfData.fSliceCandidates, sizeInBytes,
0454                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0455
0456     // Allocate and copy candidates lists
0457     fSurfDataStaging.fNStates = surfData.fNStates;
0458     sizeInBytes               = sizeof(surfData.fCandidates[0]) * surfData.fNStates;
0459     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fCandidates, sizeInBytes));
0460     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fCandidates, surfData.fCandidates, sizeInBytes,
0461                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0462
0463     // Nota bene: the fCandidates[i].{fCandidates,fFrameInd,fSides} are backed by the
0464     // following array and set via BrepCudaManagerFinishTransfer.
0465     fSurfDataStaging.fSizeCandList = surfData.fSizeCandList;
0466     sizeInBytes                    = sizeof(surfData.fCandList[0]) * surfData.fSizeCandList;
0467     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfDataStaging.fCandList, sizeInBytes));
0468     VECGEOM_DEVICE_API_CALL(Memcpy(fSurfDataStaging.fCandList, surfData.fCandList, sizeInBytes,
0469                                    VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0470
0471     // Now copy the staged data to the GPU
0472     VECGEOM_DEVICE_API_CALL(Malloc(&fSurfData, sizeof(SurfData_t)));
0473     VECGEOM_DEVICE_API_CALL(
0474         Memcpy(fSurfData, &fSurfDataStaging, sizeof(SurfData_t), VECGEOM_DEVICE_API_SYMBOL(MemcpyHostToDevice)));
0475
0476     // Finally finish the transfer by calling a kernel to write some pointers
0477     BrepCudaManagerFinishTransfer<<<1, 1>>>(fSurfData);
0478     VECGEOM_DEVICE_API_CALL(DeviceSynchronize());
0479   }
0480
0481   void Cleanup()
0482   {
0483     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fPVolTrans));
0484     fSurfDataStaging.fPVolTrans = nullptr;
0485     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fEllipData));
0486     fSurfDataStaging.fEllipData = nullptr;
0487     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fTorusData));
0488     fSurfDataStaging.fTorusData = nullptr;
0489     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fArb4Data));
0490     fSurfDataStaging.fArb4Data = nullptr;
0491     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fShells));
0492     fSurfDataStaging.fShells = nullptr;
0493     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fSurfShellList));
0494     fSurfDataStaging.fSurfShellList = nullptr;
0495     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fLogicList));
0496     fSurfDataStaging.fLogicList = nullptr;
0497     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fShellExitingSurfaceList));
0498     fSurfDataStaging.fShellExitingSurfaceList = nullptr;
0499     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fShellEnteringSurfaceList));
0500     fSurfDataStaging.fShellEnteringSurfaceList = nullptr;
0501     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fShellEnteringSurfacePvolList));
0502     fSurfDataStaging.fShellEnteringSurfacePvolList = nullptr;
0503     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fShellEnteringSurfacePvolTransList));
0504     fSurfDataStaging.fShellEnteringSurfacePvolTransList = nullptr;
0505     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fLocalSurf));
0506     fSurfDataStaging.fLocalSurf = nullptr;
0507     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fFramedSurf));
0508     fSurfDataStaging.fFramedSurf = nullptr;
0509     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fWindowMasks));
0510     fSurfDataStaging.fWindowMasks = nullptr;
0511     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fRingMasks));
0512     fSurfDataStaging.fRingMasks = nullptr;
0513     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fZPhiMasks));
0514     fSurfDataStaging.fZPhiMasks = nullptr;
0515     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fQuadMasks));
0516     fSurfDataStaging.fQuadMasks = nullptr;
0517     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fCommonSurfaces));
0518     fSurfDataStaging.fCommonSurfaces = nullptr;
0519     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fSides));
0520     fSurfDataStaging.fSides = nullptr;
0521     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fCandidates));
0522     fSurfDataStaging.fCandidates = nullptr;
0523     VECGEOM_DEVICE_API_CALL(Free(fSurfDataStaging.fCandList));
0524     fSurfDataStaging.fCandList = nullptr;
0525     VECGEOM_DEVICE_API_CALL(Free(fSurfData));
0526     fSurfData = nullptr;
0527   }
0528 };
0529
0530 } // namespace vgbrep
0531
0532 #endif