Gammas split

apt-sim · Sep 24, 2024 · 0170e64 · 0170e64
1 parent 5ea9d08
commit 0170e64
Show file tree

Hide file tree

Showing 6 changed files with 439 additions and 6 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -139,7 +139,7 @@ add_compile_options("$<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CONFIG:RelWithDebInfo>>:
 # - For Debug, generate full debug information - this completely disables optimizations!
 add_compile_options("$<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CONFIG:Debug>>:--device-debug>")
 # - For both, interleave the source in PTX to enhance the debugging experience.
-add_compile_options("$<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<OR:$<CONFIG:RelWithDebInfo>,$<CONFIG:Debug>>>:--source-in-ptx>")
+add_compile_options("$<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<OR:$<CONFIG:RelWithDebInfo>,$<CONFIG:Debug>>>:-G>")
 
 # Disable warnings from the CUDA frontend about unknown GCC pragmas - let the compiler decide what it likes.
 add_compile_options("$<$<COMPILE_LANGUAGE:CUDA>:-Xcudafe;--diag_suppress=unrecognized_gcc_pragma>")
@@ -153,6 +153,8 @@ if(NOT WITH_FLUCT)
   add_compile_definitions(NOFLUCTUATION)
 endif()
 
+# string(APPEND CMAKE_CUDA_FLAGS " -Xptxas=-v") 
+
 #----------------------------------------------------------------------------#
 # Build Targets
 #----------------------------------------------------------------------------#

diff --git a/examples/Example1/CMakeLists.txt b/examples/Example1/CMakeLists.txt
@@ -74,6 +74,8 @@ configure_file("macros/example1_ttbar_noadept.mac.in" "${PROJECT_BINARY_DIR}/exa
 
 # Tests
 
+string(APPEND CMAKE_CUDA_FLAGS " -Xptxas=-v") 
+
 add_test(NAME example1
   COMMAND $<TARGET_FILE:example1> -m ${PROJECT_BINARY_DIR}/example1_large_stack.mac
 )
diff --git a/examples/Example1/macros/example1_ttbar_LHCb.mac.in b/examples/Example1/macros/example1_ttbar_LHCb.mac.in
@@ -7,7 +7,7 @@
 ## Geant4 macro for modelling simplified sampling calorimeters
 ## =============================================================================
 ##
-/run/numberOfThreads 1
+/run/numberOfThreads 8
 /control/verbose 0
 /run/verbose 0
 /process/verbose 0
@@ -26,7 +26,7 @@
 /adept/setMillionsOfTrackSlots 8
 /adept/setMillionsOfHitSlots 1
 ## Device stack limit
-# /adept/setCUDAStackLimit 4096
+/adept/setCUDAStackLimit 4096
 
 # If true, particles are transported on the GPU across the whole geometry, GPU regions are ignored
 /adept/setTrackInAllRegions true

diff --git a/include/AdePT/core/AdePTTransport.cuh b/include/AdePT/core/AdePTTransport.cuh
@@ -10,7 +10,11 @@
 #include <AdePT/navigation/AdePTNavigator.h>
 #include <AdePT/base/MParray.h>
 #include <AdePT/kernels/electrons.cuh>
-#include <AdePT/kernels/gammas.cuh>
+// #include <AdePT/kernels/gammas.cuh>
+
+
+#include <AdePT/kernels/gammas_experimental.cuh>
+
 
 #include <VecGeom/base/Config.h>
 #ifdef VECGEOM_ENABLE_CUDA
@@ -407,8 +411,21 @@ void ShowerGPU(IntegrationLayer &integration, int event, adeptint::TrackBuffer &
       transportBlocks = (numGammas + TransportThreads - 1) / TransportThreads;
       transportBlocks = std::min(transportBlocks, MaxBlocks);
 #endif
-      TransportGammas<AdeptScoring><<<transportBlocks, TransportThreads, 0, gammas.stream>>>(
-          gammas.trackmgr, secondaries, gammas.leakedTracks, scoring_dev, VolAuxArray::GetInstance().fAuxData_dev);
+      Physics1<<<transportBlocks, TransportThreads, 0, gammas.stream>>>(
+        gammas.trackmgr, VolAuxArray::GetInstance().fAuxData_dev
+      );
+      Transport1<<<transportBlocks, TransportThreads, 0, gammas.stream>>>(
+        gammas.trackmgr, VolAuxArray::GetInstance().fAuxData_dev
+      );
+      Relocation<<<transportBlocks, TransportThreads, 0, gammas.stream>>>(
+        gammas.trackmgr, gammas.leakedTracks, VolAuxArray::GetInstance().fAuxData_dev
+      );
+      Physics2<AdeptScoring><<<transportBlocks, TransportThreads, 0, gammas.stream>>>(
+        gammas.trackmgr, secondaries, scoring_dev, VolAuxArray::GetInstance().fAuxData_dev
+      );
+
+      // TransportGammas<AdeptScoring><<<transportBlocks, TransportThreads, 0, gammas.stream>>>(
+      //     gammas.trackmgr, secondaries, gammas.leakedTracks, scoring_dev, VolAuxArray::GetInstance().fAuxData_dev);
 
       COPCORE_CUDA_CHECK(cudaEventRecord(gammas.event, gammas.stream));
       COPCORE_CUDA_CHECK(cudaStreamWaitEvent(gpuState.stream, gammas.event, 0));

diff --git a/include/AdePT/core/Track.cuh b/include/AdePT/core/Track.cuh
@@ -17,9 +17,11 @@ struct Track {
   using Precision = vecgeom::Precision;
 
   int parentID{0}; // Stores the track id of the initial particle given to AdePT
+  int hitsurfID{0};
 
   RanluxppDouble rngState;
   double eKin;
+  double preStepEKin;
   double numIALeft[3];
   double initialRange;
   double dynamicRangeFactor;
@@ -30,8 +32,22 @@ struct Track {
   double properTime{0};
 
   vecgeom::Vector3D<Precision> pos;
+  vecgeom::Vector3D<Precision> preStepPos;
   vecgeom::Vector3D<Precision> dir;
+  vecgeom::Vector3D<Precision> preStepDir;
   vecgeom::NavigationState navState;
+  vecgeom::NavigationState nextState;
+  vecgeom::NavigationState preStepNavState;
+
+  // Variables used to store navigation results
+  bool reachedInteractionPoint{false};
+  double geometryStepLength{0};
+
+  // Variables used to store physics results from G4HepEM
+  double geometricalStepLengthFromPhysics{0};
+  int winnerProcessIndex{0};
+  double preStepMFPs[3];
+  double PEmxSec{0};
 
   __host__ __device__ double Uniform() { return rngState.Rndm(); }