Skip to content

Commit

Permalink
[async] Adapt AsyncExample after rebase on master.
Browse files Browse the repository at this point in the history
  • Loading branch information
@hageboeck
hageboeck committed Nov 14, 2024
1 parent b6c5fdd commit 61d9da8
Show file tree
Hide file tree
Showing 13 changed files with 133 additions and 372,552 deletions.
6 changes: 3 additions & 3 deletions examples/AsyncExample/AdeptIntegration.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -33,10 +33,10 @@
std::shared_ptr<AdePTTransportInterface> AdePTTransportFactory(unsigned int nThread, unsigned int nTrackSlot,
unsigned int nHitSlot, int verbosity,
std::vector<std::string> const *GPURegionNames,
bool trackInAllRegions)
bool trackInAllRegions, int cudaStackSize)
{
static std::shared_ptr<AsyncAdePT::AdeptIntegration> adePT{
new AsyncAdePT::AdeptIntegration(nThread, nTrackSlot, nHitSlot, verbosity, GPURegionNames, trackInAllRegions)};
static std::shared_ptr<AsyncAdePT::AdeptIntegration> adePT{new AsyncAdePT::AdeptIntegration(
nThread, nTrackSlot, nHitSlot, verbosity, GPURegionNames, trackInAllRegions, cudaStackSize)};
return adePT;
}

Expand Down
15 changes: 7 additions & 8 deletions examples/AsyncExample/AdeptIntegration.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -107,7 +107,7 @@ TrackBuffer::TrackBuffer(unsigned int numToDevice, unsigned int numFromDevice, u

AdeptIntegration::AdeptIntegration(unsigned short nThread, unsigned int trackCapacity, unsigned int hitBufferCapacity,
int debugLevel, std::vector<std::string> const *GPURegionNames,
bool trackInAllRegions)
bool trackInAllRegions, int cudaStackSize)
: fNThread{nThread}, fTrackCapacity{trackCapacity}, fScoringCapacity{hitBufferCapacity}, fDebugLevel{debugLevel},
fG4Integrations(nThread), fEventStates(nThread), fGPUNetEnergy(nThread, 0.),
fTrackInAllRegions{trackInAllRegions}, fGPURegionNames{GPURegionNames}
Expand All @@ -119,6 +119,11 @@ AdeptIntegration::AdeptIntegration(unsigned short nThread, unsigned int trackCap
std::atomic_init(&eventState, EventState::ScoringRetrieved);
}

if (cudaStackSize > 0) {
G4cout << "Instantiate AdePT with cuda stack size " << cudaStackSize << "\n";
COPCORE_CUDA_CHECK(cudaDeviceSetLimit(cudaLimitStackSize, cudaStackSize));
}

AdeptIntegration::FullInit();
}

Expand Down Expand Up @@ -476,12 +481,6 @@ __global__ void InitSlotManagers(SlotManager *mgr, std::size_t N)

bool AdeptIntegration::InitializeGeometry(const vecgeom::cxx::VPlacedVolume *world)
{
#ifndef NDEBUG
COPCORE_CUDA_CHECK(vecgeom::cxx::CudaDeviceSetStackLimit(16384 * 2));
#else
COPCORE_CUDA_CHECK(vecgeom::cxx::CudaDeviceSetStackLimit(16384));
#endif

// Upload geometry to GPU.
auto &cudaManager = vecgeom::cxx::CudaManager::Instance();
cudaManager.LoadGeometry(world);
Expand Down Expand Up @@ -1130,4 +1129,4 @@ void AdeptIntegration::HitProcessingLoop(HitProcessingContext *const context)
fCV_G4Workers.notify_all();
}
}
}
}
3 changes: 2 additions & 1 deletion examples/AsyncExample/AdeptIntegration.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -89,7 +89,7 @@ class AdeptIntegration : public AdePTTransportInterface {

public:
AdeptIntegration(unsigned short nThread, unsigned int trackCapacity, unsigned int hitBufferCapacity, int debugLevel,
std::vector<std::string> const *GPURegionNames, bool trackInAllRegions);
std::vector<std::string> const *GPURegionNames, bool trackInAllRegions, int cudaStackSize);
AdeptIntegration(const AdeptIntegration &other) = delete;
~AdeptIntegration();

Expand All @@ -110,6 +110,7 @@ class AdeptIntegration : public AdePTTransportInterface {
void SetTrackInAllRegions(bool trackInAllRegions) override { fTrackInAllRegions = trackInAllRegions; }
bool GetTrackInAllRegions() const override { return fTrackInAllRegions; }
void SetGPURegionNames(std::vector<std::string> const *regionNames) override { fGPURegionNames = regionNames; }
void SetCUDAStackLimit(int limit) override{};
std::vector<std::string> const *GetGPURegionNames() override { return fGPURegionNames; }
/// No effect
void Initialize(bool) override {}
Expand Down
4 changes: 2 additions & 2 deletions examples/AsyncExample/BasicScoring.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -224,14 +224,14 @@ __device__ void RecordHit(AsyncAdePT::PerEventScoring * /*scoring*/, int aParent
aGPUHit.fStepLength = aStepLength;
aGPUHit.fTotalEnergyDeposit = aTotalEnergyDeposit;
// Pre step point
aGPUHit.fPreStepPoint.fNavigationStateIndex = aPreState->GetNavIndex();
aGPUHit.fPreStepPoint.fNavigationState = *aPreState;
Copy3DVector(*aPrePosition, aGPUHit.fPreStepPoint.fPosition);
Copy3DVector(*aPreMomentumDirection, aGPUHit.fPreStepPoint.fMomentumDirection);
// Copy3DVector(aPrePolarization, aGPUHit.fPreStepPoint.fPolarization);
aGPUHit.fPreStepPoint.fEKin = aPreEKin;
aGPUHit.fPreStepPoint.fCharge = aPreCharge;
// Post step point
aGPUHit.fPostStepPoint.fNavigationStateIndex = aPostState->GetNavIndex();
aGPUHit.fPostStepPoint.fNavigationState = *aPostState;
Copy3DVector(*aPostPosition, aGPUHit.fPostStepPoint.fPosition);
Copy3DVector(*aPostMomentumDirection, aGPUHit.fPostStepPoint.fMomentumDirection);
// Copy3DVector(aPostPolarization, aGPUHit.fPostStepPoint.fPolarization);
Expand Down
31 changes: 22 additions & 9 deletions examples/AsyncExample/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -59,21 +59,32 @@ if(HepMC3_FOUND)
target_link_libraries(ExamplesCommon PRIVATE HepMC3::HepMC3)
endif()

add_library(AdePTAsyncTransport SHARED
# Include the C++ side of the AdePT-G4 integration from core AdePT:
list(TRANSFORM ADEPT_G4_INTEGRATION_SRCS PREPEND ../../)

add_library(AdePT_G4_integration_async SHARED
${ADEPT_G4_INTEGRATION_SRCS}
AdeptIntegration.cpp
AdeptIntegration.cu
BasicScoring.cu
ResourceManagement.cu)
target_link_libraries(AdePTAsyncTransport
target_link_libraries(AdePT_G4_integration_async
PUBLIC
CopCore
VecGeom::vgdml
VecGeom::vecgeomcuda_static
G4HepEm::g4HepEm
)
target_include_directories(AdePT_G4_integration_async
PUBLIC
AdePT_G4_integration
VecGeom::vecgeomcuda_static
)
set_target_properties(AdePTAsyncTransport PROPERTIES
$<BUILD_INTERFACE:${CMAKE_SOURCE_DIR}/include>
$<INSTALL_INTERFACE:${CMAKE_INSTALL_INCLUDEDIR}/>
)
set_target_properties(AdePT_G4_integration_async PROPERTIES
CUDA_SEPARABLE_COMPILATION ON
CUDA_RESOLVE_DEVICE_SYMBOLS ON)

option(AsyncTransport "Link to AdePTAsyncTransport" ON)
option(AsyncTransport "Link to AdePT_G4_integration_async" ON)

find_package(ROOT REQUIRED COMPONENTS Hist RIO)

Expand All @@ -94,12 +105,14 @@ target_link_libraries(AsyncExample
)

if(AsyncTransport)
target_link_libraries(AsyncExample PRIVATE AdePTAsyncTransport)
target_link_libraries(AsyncExample PRIVATE AdePT_G4_integration_async)
else()
target_link_libraries(AsyncExample PRIVATE AdePTTransport)
target_link_libraries(AsyncExample PRIVATE AdePT_G4_integration)
endif()

target_compile_options(AsyncExample PRIVATE -Wall -Wextra)

# Install macros and geometry file
SET(GDML ${PROJECT_BINARY_DIR}/cms2018_sd.gdml)
configure_file("macros/AsyncExample.mac.in" "${CMAKE_BINARY_DIR}/AsyncExample.mac")
configure_file("macros/AsyncExample_ttbar.mac.in" "${CMAKE_BINARY_DIR}/AsyncExample_ttbar.mac")
4 changes: 3 additions & 1 deletion examples/AsyncExample/electrons.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -189,11 +189,13 @@ static __device__ __forceinline__ void TransportElectrons(Track *electrons, cons

// Check if there's a volume boundary in between.
bool propagated = true;
long hitsurf_index = -1;
double geometryStepLength;
vecgeom::NavStateIndex nextState;
if (BzFieldValue != 0) {
geometryStepLength = fieldPropagatorBz.ComputeStepAndNextVolume<BVHNavigator>(
eKin, restMass, Charge, geometricalStepLengthFromPhysics, pos, dir, navState, nextState, propagated, safety);
eKin, restMass, Charge, geometricalStepLengthFromPhysics, pos, dir, navState, nextState, hitsurf_index,
propagated, safety);
} else {
geometryStepLength = BVHNavigator::ComputeStepAndNextVolume(pos, dir, geometricalStepLengthFromPhysics, navState,
nextState, kPush);
Expand Down
144 changes: 29 additions & 115 deletions examples/AsyncExample/macros/AsyncExample.mac.in
View file
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
# SPDX-FileCopyrightText: 2022 CERN
# SPDX-License-Identifier: Apache-2.0
# adeptint.in
# AsyncExample.in
#

## =============================================================================
Expand All @@ -12,147 +12,61 @@
/run/verbose 0
/process/verbose 0
/tracking/verbose 0
##
/adeptint/detector/filename @CMS2018_GDML@
/adeptint/detector/regionname EcalRegion
/adeptint/adept/verbose 1
## Threshold for buffering tracks before sending to GPU
/adeptint/adept/threshold 200
## Total number of GPU track slots (not per thread)
/adeptint/adept/milliontrackslots 20

/adeptint/detector/addsensitivevolume EAPD_01
/adeptint/detector/addsensitivevolume EAPD_02
/adeptint/detector/addsensitivevolume EAPD_03
/adeptint/detector/addsensitivevolume EAPD_04
/adeptint/detector/addsensitivevolume EAPD_05
/adeptint/detector/addsensitivevolume EAPD_06
/adeptint/detector/addsensitivevolume EAPD_07
/adeptint/detector/addsensitivevolume EAPD_08
/adeptint/detector/addsensitivevolume EAPD_09
/adeptint/detector/addsensitivevolume EAPD_10
/adeptint/detector/addsensitivevolume EAPD_11
/adeptint/detector/addsensitivevolume EAPD_12
/adeptint/detector/addsensitivevolume EAPD_13
/adeptint/detector/addsensitivevolume EAPD_14
/adeptint/detector/addsensitivevolume EAPD_15
/adeptint/detector/addsensitivevolume EAPD_16
/adeptint/detector/addsensitivevolume EAPD_17

/adeptint/detector/sensitivegroup EAPD

/adeptint/detector/addsensitivevolume EBRY_01
/adeptint/detector/addsensitivevolume EBRY_02
/adeptint/detector/addsensitivevolume EBRY_03
/adeptint/detector/addsensitivevolume EBRY_04
/adeptint/detector/addsensitivevolume EBRY_05
/adeptint/detector/addsensitivevolume EBRY_06
/adeptint/detector/addsensitivevolume EBRY_07
/adeptint/detector/addsensitivevolume EBRY_08
/adeptint/detector/addsensitivevolume EBRY_09
/adeptint/detector/addsensitivevolume EBRY_10
/adeptint/detector/addsensitivevolume EBRY_11
/adeptint/detector/addsensitivevolume EBRY_12
/adeptint/detector/addsensitivevolume EBRY_13
/adeptint/detector/addsensitivevolume EBRY_14
/adeptint/detector/addsensitivevolume EBRY_15
/adeptint/detector/addsensitivevolume EBRY_16
/adeptint/detector/addsensitivevolume EBRY_17
/event/verbose 0

/adeptint/detector/sensitivegroup EBRY_01
/adeptint/detector/sensitivegroup EBRY_02
/adeptint/detector/sensitivegroup EBRY_03
/adeptint/detector/sensitivegroup EBRY_04
/adeptint/detector/sensitivegroup EBRY_05
/adeptint/detector/sensitivegroup EBRY_06
/adeptint/detector/sensitivegroup EBRY_07
/adeptint/detector/sensitivegroup EBRY_08
/adeptint/detector/sensitivegroup EBRY_09
/adeptint/detector/sensitivegroup EBRY_10
/adeptint/detector/sensitivegroup EBRY_11
/adeptint/detector/sensitivegroup EBRY_12
/adeptint/detector/sensitivegroup EBRY_13
/adeptint/detector/sensitivegroup EBRY_14
/adeptint/detector/sensitivegroup EBRY_15
/adeptint/detector/sensitivegroup EBRY_16
/adeptint/detector/sensitivegroup EBRY_17

/adeptint/detector/addsensitivevolume EATJ_01
/adeptint/detector/addsensitivevolume EATJ_02
/adeptint/detector/addsensitivevolume EATJ_03
/adeptint/detector/addsensitivevolume EATJ_04
/adeptint/detector/addsensitivevolume EATJ_05
/adeptint/detector/addsensitivevolume EATJ_06
/adeptint/detector/addsensitivevolume EATJ_07
/adeptint/detector/addsensitivevolume EATJ_08
/adeptint/detector/addsensitivevolume EATJ_09
/adeptint/detector/addsensitivevolume EATJ_10
/adeptint/detector/addsensitivevolume EATJ_11
/adeptint/detector/addsensitivevolume EATJ_12
/adeptint/detector/addsensitivevolume EATJ_13
/adeptint/detector/addsensitivevolume EATJ_14
/adeptint/detector/addsensitivevolume EATJ_15
/adeptint/detector/addsensitivevolume EATJ_16
/adeptint/detector/addsensitivevolume EATJ_17

/adeptint/detector/sensitivegroup EATJ
/detector/filename @GDML@
# Temporary workaround since we don't have a G4 to VecGeom converter
/adept/setVecGeomGDML @GDML@
/adept/setVerbosity 1
## Total number of GPU track slots (not per thread)
/adept/setMillionsOfTrackSlots 16
/adept/setMillionsOfHitSlots 1
# /adept/setCUDAStackLimit 16384

/adeptint/detector/addsensitivevolume EFRY
# If true, particles are transported on the GPU across the whole geometry, GPU regions are ignored
/adept/setTrackInAllRegions true
# In order to do the GPU transport only in specific regions
/adept/addGPURegion EcalRegion
/adept/addGPURegion HcalRegion

/adeptint/detector/sensitivegroup EFRY

## -----------------------------------------------------------------------------
## Optionally, set a constant magnetic filed:
## -----------------------------------------------------------------------------
#/adeptint/detector/setField 0 0 3.8 tesla
/adeptint/detector/setField 0 0 0 tesla
/detector/setField 0 0 0 tesla
#/detector/setField 0 0 3.8 tesla

##
## -----------------------------------------------------------------------------
## Set the physics list (more exactly, the EM physics constructor):
## = 'HepEm' : the G4HepEm EM physics c.t.r.
## = 'G4Em' : the G4 EM physics c.t.r. that corresponds to G4HepEm
## = 'emstandard_opt0' : the original, G4 EM-Opt0 physics c.t.r.
## -----------------------------------------------------------------------------
##/testem/phys/addPhysics HepEm
##/testem/phys/addPhysics G4Em
##
## -----------------------------------------------------------------------------
## Set secondary production threshold, init. the run and set primary properties
## -----------------------------------------------------------------------------
/run/setCut 0.7 mm

# Do or don't initialise AdePT (consumes GPU memory)
/adeptint/adept/activate true

/run/initialize

## Event verbosity: 1 = total edep, 2 = energy deposit per placed sensitive volume
/adeptint/event/verbose 2
## User-defined Event verbosity: 1 = total edep, 2 = energy deposit per placed sensitive volume
/eventAction/verbose 2

/gun/setDefault
/gun/particle e-
/gun/energy 10 GeV
/gun/number 200
/gun/position 0 0 0
/gun/randomizeGun false
/gun/print true

# If false, the following parameters are ignored
/gun/randomizeGun true
# Usage: /gun/addParticle type ["weight" weight] ["energy" energy unit]
/gun/addParticle e- weight 1 energy 10 GeV
/gun/print false
/gun/addParticle proton weight 0 energy 10 GeV
/gun/minPhi 0 deg
/gun/maxPhi 360 deg
/gun/minTheta 10 deg
/gun/maxTheta 170 deg

##
## -----------------------------------------------------------------------------
## Run the simulation with the given number of events and print list of processes
## -----------------------------------------------------------------------------
##/tracking/verbose 1
##/process/list

# run events with parametrised simulation
# by default all created models are active
/param/ActivateModel AdePT
#/analysis/setFileName adeptint_fastsim_100events.root
/run/beamOn 4

# run events with full (detailed) simulation
/param/InActivateModel AdePT
#/analysis/setFileName adeptint_fullsim_100events.root
/run/beamOn 4
Loading

0 comments on commit 61d9da8

Please sign in to comment.