Skip to content

Commit

Permalink
First version of classes for Runge-Kutta integration of track in field
Browse files Browse the repository at this point in the history
- MagneticFieldEquation
- DormandPrince45 - Stepper
- RkIntegrationDriver
- fieldConstants.h with kB2C, delta_intersection parameters
- fieldPropagatorRungeKutta : new class for use by electrons.cu, used in new Example14

More:

Example15: field propagation using Runge-Kutta.  Based on Example13 otherwise.

Tests - unit test test_magfieldRK.cpp
  Simple checks for equation, stepper and driver classes
  Driver: check version of Advance and V1 (old)
  Checks driver vs helix results (on cpu).

Details:
--------
RkIntegrationDriver
   Introduced new Advance method.  Using AdvanceV1 as backward compatible.
   Enabled used of RK classes with double integrands (field remains float.)

fieldPropagatorConstBz:  boundary cross location accurate to within delta_intersection

 PrintFieldVectors: auxiliary methods, initially host-only
 Changed VECCORE_ATT_HOST_DEVICE to __host__ __device__

example15: Added ability to print Track info
           check result of RK integration in electrons.cu
           report differences
           Optional argument for Bz field value.
           Use TrackML as default geometry.
  • Loading branch information
jonapost committed Feb 22, 2022
1 parent 7a9d562 commit 61d6bd8
Show file tree
Hide file tree
Showing 28 changed files with 4,333 additions and 88 deletions.
1 change: 1 addition & 0 deletions examples/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -16,6 +16,7 @@ add_subdirectory(Example10)
add_subdirectory(Example11)
add_subdirectory(Example12)
add_subdirectory(Example13)
add_subdirectory(Example15)
add_subdirectory(TestEm3)
add_subdirectory(TestEm3MT)
add_subdirectory(ECS)
2 changes: 1 addition & 1 deletion examples/Example11/README.md
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,7 @@ Results are reproducible using one RANLUX++ state per track.
### Kernels

This example uses one stream per particle type to launch kernels asynchronously.
They are synchronized via a forth stream using CUDA events.
They are synchronized via a fourth stream using CUDA events.

#### `TransportElectrons<bool IsElectron>`

Expand Down
2 changes: 1 addition & 1 deletion examples/Example12/README.md
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@ Additionally, the kernels score energy deposit and the charged track length per
### Kernels

This example uses one stream per particle type to launch kernels asynchronously.
They are synchronized via a forth stream using CUDA events.
They are synchronized via a fourth stream using CUDA events.

#### `TransportElectrons<bool IsElectron>`

Expand Down
2 changes: 1 addition & 1 deletion examples/Example13/README.md
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@ Additionally, the kernels score energy deposit and the charged track length per
### Kernels

This example uses one stream per particle type to launch kernels asynchronously.
They are synchronized via a forth stream using CUDA events.
They are synchronized via a fourth stream using CUDA events.

#### `TransportElectrons<bool IsElectron>`

Expand Down
42 changes: 42 additions & 0 deletions examples/Example15/CMakeLists.txt
Original file line number Diff line number Diff line change
@@ -0,0 +1,42 @@
# SPDX-FileCopyrightText: 2021 CERN
# SPDX-License-Identifier: Apache-2.0

if(NOT TARGET G4HepEm::g4HepEm)
message(STATUS "Disabling example15 (needs G4HepEm)")
return()
endif()

if(Geant4_FOUND)
if(NOT Geant4_gdml_FOUND)
message(STATUS "Disabling example15 (needs Geant4 with GDML support)")
return()
endif()
else()
message(STATUS "Disabling example15 (needs Geant4)")
return()
endif()

# example15 is the AdePT demo example using material cuts as defined in the input gdml file
add_executable(example15
example15.cpp
example15.cu
electrons.cu
gammas.cu)
target_link_libraries(example15
PRIVATE
AdePT
CopCore::CopCore
VecGeom::vecgeom
VecGeom::vecgeomcuda_static
VecGeom::vgdml
${Geant4_LIBRARIES}
G4HepEm::g4HepEmData
G4HepEm::g4HepEmInit
G4HepEm::g4HepEmRun
CUDA::cudart)

set_target_properties(example15 PROPERTIES CUDA_SEPARABLE_COMPILATION ON CUDA_RESOLVE_DEVICE_SYMBOLS ON)

# Tests
add_test(NAME example15
COMMAND $<TARGET_FILE:example15> -gdml_file ${CMAKE_BINARY_DIR}/cms2018.gdml)
64 changes: 64 additions & 0 deletions examples/Example15/README.md
Original file line number Diff line number Diff line change
@@ -0,0 +1,64 @@
<!--
SPDX-FileCopyrightText: 2021 CERN
SPDX-License-Identifier: CC-BY-4.0
-->

## Example 14

Example demonstrating particle transportation on GPUs with a non-uniform magnetic field - in arbitrary geometry read from a GDML file.

New feature is
- definition of a non-uniform magnetic field
- integration of the equation of motion of charged particles using embedded Runge-Kutta method, Dormand Prince 4(5).
Caveats (at present) :
- fixed accuracy of integration (constant in source)

The features carried over from Example13 are:
* arbitrary geometry via gdml file (tested with cms2018.gdml from VecGeom persistency/gdml/gdmls folder) and optionally a magnetic field with constant Bz,
* geometry read as Geant4 geometry, reading in regions and cuts, to initialize G4HepEm data
* geometry read then into VecGeom, and synchronized to GPU
* G4HepEm material-cuts couple indices mapped to VecGeom logical volume id's
* physics processes for e-/e+ (including MSC) and gammas using G4HepEm.
* scoring per placed volume, no sensitive detector feature

Electrons, positrons, and gammas are stored in separate containers in device memory.
Free positions in the storage are handed out with monotonic slot numbers, slots are not reused.
Active tracks are passed via three queues per particle type (see `struct ParticleQueues`).
Results are reproducible using one RANLUX++ state per track.

Additionally, the kernels score energy deposit and the charged track length per volume.

### Kernels

This example uses one stream per particle type to launch kernels asynchronously.
They are synchronized via a fourth stream using CUDA events.

#### `TransportElectrons<bool IsElectron>`

1. Obtain safety unless the track is currently on a boundary.
2. Determine physics step limit, including conversion to geometric step length according to MSC.
3. Query geometry (or optionally magnetic field) to get geometry step length.
4. Convert geometry to true step length according to MSC, apply net direction change and discplacement.
5. Apply continuous effects; kill track if stopped.
6. If the particle reaches a boundary, perform relocation.
7. If not, and if there is a discrete process:
1. Sample the final state.
2. Update the primary and produce secondaries.

#### `TransportGammas`

1. Determine the physics step limit.
2. Query VecGeom to get geometry step length (no magnetic field for neutral particles!).
3. If the particle reaches a boundary, perform relocation.
4. If not, and if there is a discrete process:
1. Sample the final state.
2. Update the primary and produce secondaries.

#### `FinishIteration`

Clear the queues and return the tracks in flight.
This kernel runs after all secondary particles were produced.

#### `InitPrimaries` and `InitParticleQueues`

Used to initialize multiple primary particles with separate seeds.
Loading

0 comments on commit 61d6bd8

Please sign in to comment.