diff --git a/Examples/Hadr04/src/run_50events.mac b/Examples/Hadr04/src/run_50events.mac
new file mode 100644
index 00000000..30b5267d
--- /dev/null
+++ b/Examples/Hadr04/src/run_50events.mac
@@ -0,0 +1,35 @@
+# $Id$
+#
+# Macro file for "Hadr04.cc"
+# (can be run in batch, without graphic)
+#
+# neutron 2 MeV; all processes
+#
+/control/verbose 2
+/run/verbose 2
+#
+/testhadr/det/setMat Water_ts
+/testhadr/det/setSize 1 m
+#
+/testhadr/phys/thermalScattering true
+/random/setSeeds 9876 54321
+#
+/run/initialize
+#
+/process/list
+#
+/gun/particle neutron
+/gun/energy 2 MeV
+#
+/analysis/setFileName Water_thermal_44
+/analysis/h1/set 1 100 0. 100. none #nb colli >1eV
+/analysis/h1/set 2 100 0. 100. cm #track len >1eV
+/analysis/h1/set 3 100 0. 5. mus #time of flight >1eV
+/analysis/h1/set 4 100 0. 1000. none #nb colli <1eV
+/analysis/h1/set 5 100 0. 500. cm #track len <1eV
+/analysis/h1/set 6 100 0. 1000. mus #time of flight <1eV
+/analysis/h1/set 7 100 0. 500. meV #energy dist <1eV
+#
+/run/printProgress 10
+#
+/run/beamOn 50
diff --git a/README.md b/README.md
index abd75e5d..a02dc738 100644
--- a/README.md
+++ b/README.md
@@ -57,11 +57,12 @@ extra `\)` before the `;\n";`. The version in the repo is fixed, but if you
downloaded marshalgen from another source, you need to fix this)
**Install GEANT-4**
+If installing on McMaster's servers, add `. /opt/rh/devtoolset-3/enable` to your bash_profile to use the newer version of gcc.
1. (Fedora only) `yum install expat-devel`
2. `mkdir /path/to/GEANT4-GPU/geant4.10.00.p02-build /path/to/GEANT4-GPU/
geant4.10.00.p02-install`
3. `cd /path/to/GEANT4-GPU/geant4.10.00.p02-build`
-4. `cmake -DGEANT4_INSTALL_DATA=ON -DGEANT4_ENABLE_CUDA=ON -DUSE_SYSTEM_EXPAT=OFF -DCMAKE_INSTALL_PREFIX=/path/to/GEANT4-GPU/geant4.10.00.p02-install /path/to/GEANT4-GPU/geant4.10.00.p02`
+4. `cmake -DGEANT4_INSTALL_DATA=ON -DGEANT4_ENABLE_CUDA=ON -DGEANT4_USE_SYSTEM_EXPAT=OFF -DCMAKE_INSTALL_PREFIX=/path/to/GEANT4-GPU/geant4.10.00.p02-install /path/to/GEANT4-GPU/geant4.10.00.p02`
IF installing on McMaster's server, you must add flag `-DCUDA_HOST_COMPILER=/usr/bin/g++`
5. `make -jN` where `N` is the number of processors on your computer
6. `make install`
diff --git a/Results/Hadr04-Feb23_12pm-gpu1-CPU.txt b/Results/Hadr04-Feb23_12pm-gpu1-CPU.txt
new file mode 100644
index 00000000..071bed9a
--- /dev/null
+++ b/Results/Hadr04-Feb23_12pm-gpu1-CPU.txt
@@ -0,0 +1,224 @@
+*************************************************************
+ Geant4 version Name: geant4-10-02 (4-December-2015)
+ Copyright : Geant4 Collaboration
+ Reference : NIM A 506 (2003), 250-303
+ WWW : http://cern.ch/geant4
+*************************************************************
+
+/run/verbose 2
+#
+/testhadr/det/setMat Water_ts
+/testhadr/det/setSize 1 m
+/run/reinitializeGeometry
+#
+/testhadr/phys/thermalScattering true
+/random/setSeeds 9876 54321
+#
+/run/initialize
+userDetector->Construct() start.
+
+ The Box is 1 m of Water_ts
+
+ Material: Water_ts density: 1.000 g/cm3 RadL: 36.084 cm Nucl.Int.Length: 75.376 cm
+ Imean: 78.000 eV
+
+ ---> Element: TS_H_of_Water (H) Z = 1.0 N = 1 A = 1.008 g/mole
+ ---> Isotope: H1 Z = 1 N = 1 A = 1.01 g/mole abundance: 99.989 %
+ ---> Isotope: H2 Z = 1 N = 2 A = 2.01 g/mole abundance: 0.011 %
+ ElmMassFraction: 11.19 % ElmAbundance 66.67 %
+
+ ---> Element: Oxygen (O) Z = 8.0 N = 16 A = 16.000 g/mole
+ ---> Isotope: O16 Z = 8 N = 16 A = 15.99 g/mole abundance: 99.757 %
+ ---> Isotope: O17 Z = 8 N = 17 A = 17.00 g/mole abundance: 0.038 %
+ ---> Isotope: O18 Z = 8 N = 18 A = 18.00 g/mole abundance: 0.205 %
+ ElmMassFraction: 88.81 % ElmAbundance 33.33 %
+
+Water_ts is registered to the default region.
+physicsList->Construct() start.
+@@@ G4ParticleHPInelasticData instantiated for particle neutron data directory variable is G4NEUTRONHPDATA pointing to /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5
+NeutronHP: /Capture file for Z = 8, A = 18 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Capture/CrossSection/8_17_Oxygen
+NeutronHP: /Elastic file for Z = 8, A = 18 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Elastic/CrossSection/8_17_Oxygen
+NeutronHP: /Inelastic file for Z = 8, A = 18 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Inelastic/CrossSection/8_17_Oxygen
+NeutronHP: /Capture file for Z = 6, A = 12 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Capture/CrossSection/6_nat_Carbon
+NeutronHP: /Elastic file for Z = 6, A = 12 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Elastic/CrossSection/6_nat_Carbon
+NeutronHP: /Inelastic file for Z = 6, A = 12 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Inelastic/CrossSection/6_nat_Carbon
+/u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5
+@@@ G4ParticleHPInelastic instantiated for particle neutron data directory variable is G4NEUTRONHPDATA pointing to /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Inelastic
+physicsList->CheckParticleList() start.
+physicsList->setCut() start.
+#
+/process/list
+ Transportation, hadElastic, neutronInelastic, nCapture
+ nFission
+#
+/gun/particle neutron
+/gun/energy 2 MeV
+#
+/analysis/setFileName Water_thermal_44
+Set file name: Water_thermal_44
+/analysis/h1/set 1 100 0. 100. none
+/analysis/h1/set 2 100 0. 100. cm
+/analysis/h1/set 3 100 0. 5. mus
+/analysis/h1/set 4 100 0. 1000. none
+/analysis/h1/set 5 100 0. 500. cm
+/analysis/h1/set 6 100 0. 1000. mus
+/analysis/h1/set 7 100 0. 500. meV
+#
+/run/printProgress 100
+#
+/run/beamOn 500
+NeutronHP: /Capture file for Z = 8, A = 18 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Capture/CrossSection/8_17_Oxygen
+NeutronHP: /Elastic file for Z = 8, A = 18 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Elastic/CrossSection/8_17_Oxygen
+NeutronHP: /Inelastic file for Z = 8, A = 18 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Inelastic/CrossSection/8_17_Oxygen
+NeutronHP: /Capture file for Z = 6, A = 12 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Capture/CrossSection/6_nat_Carbon
+NeutronHP: /Elastic file for Z = 6, A = 12 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Elastic/CrossSection/6_nat_Carbon
+NeutronHP: /Inelastic file for Z = 6, A = 12 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Inelastic/CrossSection/6_nat_Carbon
+
+Neutron HP Thermal Scattering Data: Following material-element pairs and/or elements are registered.
+Element TS_H_of_Water, internal thermal scattering id 0
+Element TS_C_of_Graphite, internal thermal scattering id 1
+
+NeutronHP: /Elastic file for Z = 8, A = 18 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Elastic/CrossSection/8_17_Oxygen
+NeutronHP: /Elastic file for Z = 6, A = 12 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Elastic/CrossSection/6_nat_Carbon
+
+Neutron HP Thermal Scattering: Following material-element pairs or elements are registered.
+Element TS_H_of_Water, internal thermal scattering id 0
+Element TS_C_of_Graphite, internal thermal scattering id 1
+
+
+Neutron HP Thermal Scattering Data: Following material-element pairs and/or elements are registered.
+Element TS_H_of_Water, internal thermal scattering id 0
+Element TS_C_of_Graphite, internal thermal scattering id 1
+
+
+====================================================================
+ HADRONIC PROCESSES SUMMARY (verbose level 1)
+
+---------------------------------------------------
+ Hadronic Processes for neutron
+
+ Process: hadElastic
+ Model: NeutronHPElastic: 4 eV ---> 20 MeV
+ Model: NeutronHPThermalScattering: 0 meV ---> 4 eV
+ Cr_sctns: NeutronHPThermalScatteringData: 0 meV ---> 4 eV
+ Cr_sctns: NeutronHPElasticXS: 0 meV ---> 20 MeV
+ Cr_sctns: GheishaElastic: 0 meV ---> 100 TeV
+
+ Process: neutronInelastic
+ Model: NeutronHPInelastic: 0 meV ---> 20 MeV
+ Cr_sctns: NeutronHPInelasticXS: 0 meV ---> 20 MeV
+ Cr_sctns: GheishaInelastic: 0 meV ---> 100 TeV
+
+ Process: nCapture
+ Model: NeutronHPCapture: 0 meV ---> 20 MeV
+ Cr_sctns: NeutronHPCaptureXS: 0 meV ---> 20 MeV
+ Cr_sctns: GheishaCaptureXS: 0 meV ---> 100 TeV
+
+ Process: nFission
+ Model: NeutronHPFission: 0 meV ---> 20 MeV
+ Cr_sctns: NeutronHPFissionXS: 0 meV ---> 20 MeV
+ Cr_sctns: GheishaFissionXS: 0 meV ---> 100 TeV
+
+================================================================
+/u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5
+@@@ G4ParticleHPInelastic instantiated for particle neutron data directory variable is G4NEUTRONHPDATA pointing to /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Inelastic
+NeutronHP: /Capture file for Z = 8, A = 18 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Capture/CrossSection/8_17_Oxygen
+NeutronHP: /Capture file for Z = 6, A = 12 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Capture/CrossSection/6_nat_Carbon
+
+Region -- -- appears in world volume
+ This region is in the mass world.
+ Root logical volume(s) : Water_ts
+ Pointers : G4VUserRegionInformation[0], G4UserLimits[0], G4FastSimulationManager[0], G4UserSteppingAction[0]
+ Materials : Water_ts
+ Production cuts : gamma 1 mm e- 1 mm e+ 1 mm proton 0 fm
+
+Region -- -- is not associated to any world.
+ Root logical volume(s) :
+ Pointers : G4VUserRegionInformation[0], G4UserLimits[0], G4FastSimulationManager[0], G4UserSteppingAction[0]
+ Materials :
+ Production cuts : gamma 1 mm e- 1 mm e+ 1 mm proton 0 fm
+
+========= Table of registered couples ==============================
+
+Index : 0 used in the geometry : Yes
+ Material : Water_ts
+ Range cuts : gamma 1 mm e- 1 mm e+ 1 mm proton 0 fm
+ Energy thresholds : gamma 2.94056 keV e- 351.877 keV e+ 342.545 keV proton 0 meV
+ Region(s) which use this couple :
+ DefaultRegionForTheWorld
+
+====================================================================
+
+Start closing geometry.
+G4GeometryManager::ReportVoxelStats -- Voxel Statistics
+
+ Total memory consumed for geometry optimisation: 0 kByte
+ Total CPU time elapsed for geometry optimisation: 0 seconds
+### Run 0 starts.
+
+--------- Ranecu engine status ---------
+ Initial seed (index) = 0
+ Current couple of seeds = 9876, 54321
+----------------------------------------
+... open Root analysis file : Water_thermal_44.root - done
+--> Event 0 starts.
+--> Event 100 starts.
+--> Event 200 starts.
+--> Event 300 starts.
+--> Event 400 starts.
+ Run terminated.
+Run Summary
+ Number of events processed : 500
+ User=14.19s Real=14.2s Sys=0s
+
+ The run is 500 neutron of 2 MeV through 50 cm of Water_ts (density: 1 g/cm3 )
+
+ Process calls frequency :
+ hadElastic= 108818 nCapture= 500
+
+ Parcours of incident neutron:
+ nb of collisions E>1*eV= 16.63 E<1*eV= 202.01 total= 218.64
+ track length E>1*eV= 19.976 cm E<1*eV= 75.695 cm total= 95.671 cm
+ time of flight E>1*eV= 753.28 ns E<1*eV= 207.82 mus total= 208.57 mus
+
+ List of generated particles:
+ O16: 1512 Emean = 43.921 keV ( 2.1992 meV --> 446.59 keV)
+ O17: 2 Emean = 202.95 eV ( 187.15 eV --> 218.76 eV )
+ O18: 1 Emean = 9.3256 eV ( 9.3256 eV --> 9.3256 eV )
+ deuteron: 498 Emean = 1.3195 keV ( 1.2775 keV --> 1.7025 keV)
+ gamma: 502 Emean = 2.2188 MeV ( 870.8 keV --> 2.2244 MeV)
+ proton: 11474 Emean = 81.412 keV ( 0.14188 meV --> 1.9993 MeV)
+... write Root file : Water_thermal_44.root - done
+
+--------- Ranecu engine status ---------
+ Initial seed (index) = 0
+ Current couple of seeds = 1392330422, 48169388
+----------------------------------------
+G4 kernel has come to Quit state.
+UserDetectorConstruction deleted.
+UserPhysicsList deleted.
+UserActionInitialization deleted.
+UserRunAction deleted.
+UserPrimaryGenerator deleted.
+RunManager is deleting RunManagerKernel.
+EventManager deleted.
+Units table cleared.
+Total navigation history collections cleaned: 4
+================== Deleting memory pools ===================
+Pool ID '20G4NavigationLevelRep', size : 0.00385 MB
+Pool ID '24G4ReferenceCountedHandleIvE', size : 0.000961 MB
+Pool ID '7G4Event', size : 0.000961 MB
+Pool ID '15G4PrimaryVertex', size : 0.000961 MB
+Pool ID '17G4PrimaryParticle', size : 0.000961 MB
+Pool ID '17G4DynamicParticle', size : 0.0154 MB
+Pool ID '7G4Track', size : 0.0308 MB
+Pool ID '18G4TouchableHistory', size : 0.000961 MB
+Pool ID '15G4CountedObjectIvE', size : 0.000961 MB
+Pool ID '17G4ReactionProduct', size : 0.000961 MB
+Number of memory pools allocated: 10; of which, static: 0
+Dynamic pools deleted: 10 / Total memory freed: 0.057 MB
+============================================================
+G4Allocator objects are deleted.
+UImanager deleted.
+StateManager deleted.
+RunManagerKernel is deleted. Good bye :)
\ No newline at end of file
diff --git a/Results/Hadr04-Feb23_12pm-gpu1-GPU.txt b/Results/Hadr04-Feb23_12pm-gpu1-GPU.txt
new file mode 100644
index 00000000..eead06fc
--- /dev/null
+++ b/Results/Hadr04-Feb23_12pm-gpu1-GPU.txt
@@ -0,0 +1,224 @@
+*************************************************************
+ Geant4 version Name: geant4-10-02 (4-December-2015)
+ Copyright : Geant4 Collaboration
+ Reference : NIM A 506 (2003), 250-303
+ WWW : http://cern.ch/geant4
+*************************************************************
+
+/run/verbose 2
+#
+/testhadr/det/setMat Water_ts
+/testhadr/det/setSize 1 m
+/run/reinitializeGeometry
+#
+/testhadr/phys/thermalScattering true
+/random/setSeeds 9876 54321
+#
+/run/initialize
+userDetector->Construct() start.
+
+ The Box is 1 m of Water_ts
+
+ Material: Water_ts density: 1.000 g/cm3 RadL: 36.084 cm Nucl.Int.Length: 75.376 cm
+ Imean: 78.000 eV
+
+ ---> Element: TS_H_of_Water (H) Z = 1.0 N = 1 A = 1.008 g/mole
+ ---> Isotope: H1 Z = 1 N = 1 A = 1.01 g/mole abundance: 99.989 %
+ ---> Isotope: H2 Z = 1 N = 2 A = 2.01 g/mole abundance: 0.011 %
+ ElmMassFraction: 11.19 % ElmAbundance 66.67 %
+
+ ---> Element: Oxygen (O) Z = 8.0 N = 16 A = 16.000 g/mole
+ ---> Isotope: O16 Z = 8 N = 16 A = 15.99 g/mole abundance: 99.757 %
+ ---> Isotope: O17 Z = 8 N = 17 A = 17.00 g/mole abundance: 0.038 %
+ ---> Isotope: O18 Z = 8 N = 18 A = 18.00 g/mole abundance: 0.205 %
+ ElmMassFraction: 88.81 % ElmAbundance 33.33 %
+
+Water_ts is registered to the default region.
+physicsList->Construct() start.
+@@@ G4ParticleHPInelasticData instantiated for particle neutron data directory variable is G4NEUTRONHPDATA pointing to /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5
+NeutronHP: /Capture file for Z = 8, A = 18 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Capture/CrossSection/8_17_Oxygen
+NeutronHP: /Elastic file for Z = 8, A = 18 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Elastic/CrossSection/8_17_Oxygen
+NeutronHP: /Inelastic file for Z = 8, A = 18 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Inelastic/CrossSection/8_17_Oxygen
+NeutronHP: /Capture file for Z = 6, A = 12 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Capture/CrossSection/6_nat_Carbon
+NeutronHP: /Elastic file for Z = 6, A = 12 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Elastic/CrossSection/6_nat_Carbon
+NeutronHP: /Inelastic file for Z = 6, A = 12 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Inelastic/CrossSection/6_nat_Carbon
+/u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5
+@@@ G4ParticleHPInelastic instantiated for particle neutron data directory variable is G4NEUTRONHPDATA pointing to /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Inelastic
+physicsList->CheckParticleList() start.
+physicsList->setCut() start.
+#
+/process/list
+ Transportation, hadElastic, neutronInelastic, nCapture
+ nFission
+#
+/gun/particle neutron
+/gun/energy 2 MeV
+#
+/analysis/setFileName Water_thermal_44
+Set file name: Water_thermal_44
+/analysis/h1/set 1 100 0. 100. none
+/analysis/h1/set 2 100 0. 100. cm
+/analysis/h1/set 3 100 0. 5. mus
+/analysis/h1/set 4 100 0. 1000. none
+/analysis/h1/set 5 100 0. 500. cm
+/analysis/h1/set 6 100 0. 1000. mus
+/analysis/h1/set 7 100 0. 500. meV
+#
+/run/printProgress 100
+#
+/run/beamOn 500
+NeutronHP: /Capture file for Z = 8, A = 18 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Capture/CrossSection/8_17_Oxygen
+NeutronHP: /Elastic file for Z = 8, A = 18 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Elastic/CrossSection/8_17_Oxygen
+NeutronHP: /Inelastic file for Z = 8, A = 18 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Inelastic/CrossSection/8_17_Oxygen
+NeutronHP: /Capture file for Z = 6, A = 12 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Capture/CrossSection/6_nat_Carbon
+NeutronHP: /Elastic file for Z = 6, A = 12 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Elastic/CrossSection/6_nat_Carbon
+NeutronHP: /Inelastic file for Z = 6, A = 12 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Inelastic/CrossSection/6_nat_Carbon
+
+Neutron HP Thermal Scattering Data: Following material-element pairs and/or elements are registered.
+Element TS_H_of_Water, internal thermal scattering id 0
+Element TS_C_of_Graphite, internal thermal scattering id 1
+
+NeutronHP: /Elastic file for Z = 8, A = 18 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Elastic/CrossSection/8_17_Oxygen
+NeutronHP: /Elastic file for Z = 6, A = 12 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Elastic/CrossSection/6_nat_Carbon
+
+Neutron HP Thermal Scattering: Following material-element pairs or elements are registered.
+Element TS_H_of_Water, internal thermal scattering id 0
+Element TS_C_of_Graphite, internal thermal scattering id 1
+
+
+Neutron HP Thermal Scattering Data: Following material-element pairs and/or elements are registered.
+Element TS_H_of_Water, internal thermal scattering id 0
+Element TS_C_of_Graphite, internal thermal scattering id 1
+
+
+====================================================================
+ HADRONIC PROCESSES SUMMARY (verbose level 1)
+
+---------------------------------------------------
+ Hadronic Processes for neutron
+
+ Process: hadElastic
+ Model: NeutronHPElastic: 4 eV ---> 20 MeV
+ Model: NeutronHPThermalScattering: 0 meV ---> 4 eV
+ Cr_sctns: NeutronHPThermalScatteringData: 0 meV ---> 4 eV
+ Cr_sctns: NeutronHPElasticXS: 0 meV ---> 20 MeV
+ Cr_sctns: GheishaElastic: 0 meV ---> 100 TeV
+
+ Process: neutronInelastic
+ Model: NeutronHPInelastic: 0 meV ---> 20 MeV
+ Cr_sctns: NeutronHPInelasticXS: 0 meV ---> 20 MeV
+ Cr_sctns: GheishaInelastic: 0 meV ---> 100 TeV
+
+ Process: nCapture
+ Model: NeutronHPCapture: 0 meV ---> 20 MeV
+ Cr_sctns: NeutronHPCaptureXS: 0 meV ---> 20 MeV
+ Cr_sctns: GheishaCaptureXS: 0 meV ---> 100 TeV
+
+ Process: nFission
+ Model: NeutronHPFission: 0 meV ---> 20 MeV
+ Cr_sctns: NeutronHPFissionXS: 0 meV ---> 20 MeV
+ Cr_sctns: GheishaFissionXS: 0 meV ---> 100 TeV
+
+================================================================
+/u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5
+@@@ G4ParticleHPInelastic instantiated for particle neutron data directory variable is G4NEUTRONHPDATA pointing to /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Inelastic
+NeutronHP: /Capture file for Z = 8, A = 18 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Capture/CrossSection/8_17_Oxygen
+NeutronHP: /Capture file for Z = 6, A = 12 is not found and NeutronHP will use /u50/dougls2/GEANT4-GPU/geant4.10.02-install/share/Geant4-10.2.0/data/G4NDL4.5/Capture/CrossSection/6_nat_Carbon
+
+Region -- -- appears in world volume
+ This region is in the mass world.
+ Root logical volume(s) : Water_ts
+ Pointers : G4VUserRegionInformation[0], G4UserLimits[0], G4FastSimulationManager[0], G4UserSteppingAction[0]
+ Materials : Water_ts
+ Production cuts : gamma 1 mm e- 1 mm e+ 1 mm proton 0 fm
+
+Region -- -- is not associated to any world.
+ Root logical volume(s) :
+ Pointers : G4VUserRegionInformation[0], G4UserLimits[0], G4FastSimulationManager[0], G4UserSteppingAction[0]
+ Materials :
+ Production cuts : gamma 1 mm e- 1 mm e+ 1 mm proton 0 fm
+
+========= Table of registered couples ==============================
+
+Index : 0 used in the geometry : Yes
+ Material : Water_ts
+ Range cuts : gamma 1 mm e- 1 mm e+ 1 mm proton 0 fm
+ Energy thresholds : gamma 2.94056 keV e- 351.877 keV e+ 342.545 keV proton 0 meV
+ Region(s) which use this couple :
+ DefaultRegionForTheWorld
+
+====================================================================
+
+Start closing geometry.
+G4GeometryManager::ReportVoxelStats -- Voxel Statistics
+
+ Total memory consumed for geometry optimisation: 0 kByte
+ Total CPU time elapsed for geometry optimisation: 0 seconds
+### Run 0 starts.
+
+--------- Ranecu engine status ---------
+ Initial seed (index) = 0
+ Current couple of seeds = 9876, 54321
+----------------------------------------
+... open Root analysis file : Water_thermal_44.root - done
+--> Event 0 starts.
+--> Event 100 starts.
+--> Event 200 starts.
+--> Event 300 starts.
+--> Event 400 starts.
+ Run terminated.
+Run Summary
+ Number of events processed : 500
+ User=50.93s Real=58.04s Sys=7.08s
+
+ The run is 500 neutron of 2 MeV through 50 cm of Water_ts (density: 1 g/cm3 )
+
+ Process calls frequency :
+ hadElastic= 108818 nCapture= 500
+
+ Parcours of incident neutron:
+ nb of collisions E>1*eV= 16.63 E<1*eV= 202.01 total= 218.64
+ track length E>1*eV= 19.976 cm E<1*eV= 75.695 cm total= 95.671 cm
+ time of flight E>1*eV= 753.28 ns E<1*eV= 207.82 mus total= 208.57 mus
+
+ List of generated particles:
+ O16: 1512 Emean = 43.921 keV ( 2.1992 meV --> 446.59 keV)
+ O17: 2 Emean = 202.95 eV ( 187.15 eV --> 218.76 eV )
+ O18: 1 Emean = 9.3256 eV ( 9.3256 eV --> 9.3256 eV )
+ deuteron: 498 Emean = 1.3195 keV ( 1.2775 keV --> 1.7025 keV)
+ gamma: 502 Emean = 2.2188 MeV ( 870.8 keV --> 2.2244 MeV)
+ proton: 11474 Emean = 81.412 keV ( 0.14188 meV --> 1.9993 MeV)
+... write Root file : Water_thermal_44.root - done
+
+--------- Ranecu engine status ---------
+ Initial seed (index) = 0
+ Current couple of seeds = 1392330422, 48169388
+----------------------------------------
+G4 kernel has come to Quit state.
+UserDetectorConstruction deleted.
+UserPhysicsList deleted.
+UserActionInitialization deleted.
+UserRunAction deleted.
+UserPrimaryGenerator deleted.
+RunManager is deleting RunManagerKernel.
+EventManager deleted.
+Units table cleared.
+Total navigation history collections cleaned: 4
+================== Deleting memory pools ===================
+Pool ID '20G4NavigationLevelRep', size : 0.00385 MB
+Pool ID '24G4ReferenceCountedHandleIvE', size : 0.000961 MB
+Pool ID '7G4Event', size : 0.000961 MB
+Pool ID '15G4PrimaryVertex', size : 0.000961 MB
+Pool ID '17G4PrimaryParticle', size : 0.000961 MB
+Pool ID '17G4DynamicParticle', size : 0.0154 MB
+Pool ID '7G4Track', size : 0.0308 MB
+Pool ID '18G4TouchableHistory', size : 0.000961 MB
+Pool ID '15G4CountedObjectIvE', size : 0.000961 MB
+Pool ID '17G4ReactionProduct', size : 0.000961 MB
+Number of memory pools allocated: 10; of which, static: 0
+Dynamic pools deleted: 10 / Total memory freed: 0.057 MB
+============================================================
+G4Allocator objects are deleted.
+UImanager deleted.
+StateManager deleted.
+RunManagerKernel is deleted. Good bye :)
\ No newline at end of file
diff --git a/geant4.10.02/source/externals/cuda/src/G4ParticleHPVector_CUDA.cu b/geant4.10.02/source/externals/cuda/src/G4ParticleHPVector_CUDA.cu
index 61bc2935..c4527d5a 100644
--- a/geant4.10.02/source/externals/cuda/src/G4ParticleHPVector_CUDA.cu
+++ b/geant4.10.02/source/externals/cuda/src/G4ParticleHPVector_CUDA.cu
@@ -35,16 +35,16 @@ __global__ void SetAllNegativeXsecToZero_CUDA(G4ParticleHPDataPoint * theData, i
}
// http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomic-functions
-__device__ double AtomicAdd_CUDA(double* address, double val) {
- unsigned long long int* address_as_ull = (unsigned long long int*)address;
- unsigned long long int old = *address_as_ull, assumed;
- do {
- assumed = old;
- old = atomicCAS(address_as_ull, assumed,
- __double_as_longlong(val + __longlong_as_double(assumed)));
- // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
- } while (assumed != old);
- return __longlong_as_double(old);
+__device__ double AtomicAdd_CUDA(double* address, double val) {
+ unsigned long long int* address_as_ull = (unsigned long long int*)address;
+ unsigned long long int old = *address_as_ull, assumed;
+ do {
+ assumed = old;
+ old = atomicCAS(address_as_ull, assumed,
+ __double_as_longlong(val + __longlong_as_double(assumed)));
+ // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
+ } while (assumed != old);
+ return __longlong_as_double(old);
}
__global__ void CopyDataPointsToBuffer_CUDA(G4ParticleHPDataPoint * fromBuffer, G4ParticleHPDataPoint * toBuffer, G4int nEntries) {
@@ -65,7 +65,7 @@ __global__ void CopyTheIntegralToBuffer_CUDA(G4double * fromBuffer, G4double * t
/***********************************************
* Constructors, Deconstructors
***********************************************/
-G4ParticleHPVector_CUDA::G4ParticleHPVector_CUDA() {
+G4ParticleHPVector_CUDA::G4ParticleHPVector_CUDA() {
PerformInitialization(20);
}
@@ -131,7 +131,7 @@ G4ParticleHPVector_CUDA::~G4ParticleHPVector_CUDA() {
void G4ParticleHPVector_CUDA::OperatorEquals(G4ParticleHPVector_CUDA * right) {
G4int i;
-
+
totalIntegral = right->totalIntegral;
nEntries = right->nEntries;
nPoints = right->nPoints;
@@ -145,11 +145,11 @@ void G4ParticleHPVector_CUDA::OperatorEquals(G4ParticleHPVector_CUDA * right) {
cudaMalloc(&d_theIntegral, nEntries * sizeof(G4double));
CopyTheIntegralToBuffer_CUDA<<>> (right->d_theIntegral, d_theIntegral, nEntries);
}
-
+
cudaMalloc(&d_theData, nPoints * sizeof(G4ParticleHPDataPoint));
CopyDataPointsToBuffer_CUDA<<>> (right->d_theData, d_theData, nEntries);
- theManager = right->theManager;
+ theManager = right->theManager;
label = right->label;
Verbose = right->Verbose;
@@ -180,7 +180,7 @@ G4double G4ParticleHPVector_CUDA::GetX(G4int i) {
return *(h_singleDoubleResult);
}
-G4double G4ParticleHPVector_CUDA::GetY(G4int i) {
+G4double G4ParticleHPVector_CUDA::GetY(G4int i) {
if (i < 0) {
i = 0;
}
@@ -244,7 +244,7 @@ void G4ParticleHPVector_CUDA::SetXsec(G4int i, G4double x) {
void G4ParticleHPVector_CUDA::Init(std::istream & aDataFile, G4double ux, G4double uy) {
G4int total;
aDataFile >> total;
-
+
if (d_theData) {
cudaFree(d_theData);
}
@@ -280,7 +280,7 @@ __global__ void SampleLinFindLastIndex_CUDA(G4double * theIntegral, int rand, in
__global__ void SampleLinGetValues(G4ParticleHPDataPoint * theData, G4double * theIntegral, G4double * d_vals, G4int i) {
// d_vals = [x1,x2,y1,y2]
switch(threadIdx.x) {
- case 0:
+ case 0:
d_vals[0] = theIntegral[i-1];
break;
case 1:
@@ -315,7 +315,7 @@ G4double G4ParticleHPVector_CUDA::SampleLin() {
int nBlocks = GetNumBlocks(nEntries);
SampleLinFindLastIndex_CUDA<<>> (d_theIntegral, randNum, d_resultIndex, nEntries);
-
+
G4int i = 0;
cudaMemcpy(&i, d_resultIndex, sizeof(G4int), cudaMemcpyDeviceToHost);
if (i != GetVectorLength() - 1) {
@@ -326,12 +326,12 @@ G4double G4ParticleHPVector_CUDA::SampleLin() {
G4double* d_vals;
cudaMalloc(&d_vals, 4*sizeof(G4double));
SampleLinGetValues<<<1, 4>>>(d_theData, d_theIntegral, d_vals, i);
-
+
G4double vals[4];
cudaMemcpy(vals, d_vals, 4*sizeof(G4double), cudaMemcpyDeviceToHost);
-
+
result = theLin.Lin(randNum, vals[0], vals[1], vals[2], vals[3]);
-
+
cudaFree(d_resultIndex);
cudaFree(d_vals);
free(vals);
@@ -361,9 +361,9 @@ __global__ void Integrate_CUDA(G4ParticleHPDataPoint * theData, G4double * sum,
double toAdd = 0;
G4InterpolationScheme aScheme = theManager.GetScheme(i);
-
+
// NOTE: cuda's log function requires compute capability >= 3.0 for double precision
- // make sure you are compiling for 3.0 (nvcc -arch sm_30)
+ // make sure you are compiling for 3.0 (nvcc -arch sm_30)
if (aScheme == LINLIN || aScheme == CLINLIN || aScheme == ULINLIN) {
toAdd += 0.5 * (y2+y1) * (x2-x1);
}
@@ -391,13 +391,14 @@ __global__ void Integrate_CUDA(G4ParticleHPDataPoint * theData, G4double * sum,
}
}
}
+
void G4ParticleHPVector_CUDA::Integrate() {
printf("\nCUDA - Integrate (nEntries: %d)", nEntries);
if (nEntries == 1) {
totalIntegral = 0;
return;
}
-
+
G4double *d_sum;
cudaMalloc(&d_sum, sizeof(G4double));
Integrate_CUDA<<<1, nEntries>>>(d_theData, d_sum, theManager);
@@ -419,7 +420,7 @@ void G4ParticleHPVector_CUDA::IntegrateAndNormalise() {
return;
}
cudaMalloc(&d_theIntegral, nEntries * sizeof(G4double));
-
+
if (nEntries == 1) {
G4double one = 1.0;
SetValueTo_CUDA<<<1,1>>> (&d_theIntegral[0], 1.0);
@@ -482,10 +483,15 @@ void G4ParticleHPVector_CUDA::Times(G4double factor) {
/******************************************
* Functions from .cc
******************************************/
-__global__ void GetXSecFirstIndex_CUDA(G4ParticleHPDataPoint * theData, G4double e, int * resultIndex, int nEntries) {
- int idx = blockDim.x * blockIdx.x + threadIdx.x;
- if (idx < nEntries && idx < *(resultIndex) && theData[idx].energy >= e) {
- atomicMin(resultIndex, idx);
+__global__ void GetXSecFirstIndex_CUDA(G4ParticleHPDataPoint * theData, G4double e, int * resultIndex, int elementsPerThread, int nEntries) {
+ int tidx = blockDim.x * blockIdx.x + threadIdx.x;
+ int start = elementsPerThread * tidx;
+ int end = (start + elementsPerThread < nEntries) ? start + elementsPerThread : nEntries - 1;
+ for (int i = start; i < end; i++) {
+ if (i < *(resultIndex) && theData[i].energy >= e) {
+ atomicMin(resultIndex, i);
+ return;
+ }
}
}
@@ -524,21 +530,21 @@ G4double G4ParticleHPVector_CUDA::GetXsec(G4double e) {
}
// TODO try having each thread iterate over X values in array
-
// look at StorkNeutronHPCSData - line 295
-
+
SetValueTo_CUDA<<<1,1>>> (d_singleIntResult, nEntries);
- int nBlocks = GetNumBlocks(nEntries);
- GetXSecFirstIndex_CUDA<<>> (d_theData, e, d_singleIntResult, nEntries);
-
+ int elementsPerThread = 4;
+ int nBlocks = GetNumBlocks(nEntries/elementsPerThread);
+ // printf("\nnEntries = [%5d] | elementsPerThread = [%5d] | nBlocks = [%5d]", nEntries, elementsPerThread, nBlocks);
+ GetXSecFirstIndex_CUDA<<>> (d_theData, e, d_singleIntResult, elementsPerThread, nEntries);
GetYForXSec_CUDA<<<1, 1>>> (d_theData, e, d_singleIntResult, d_res, nEntries);
cudaMemcpy(h_res, d_res, sizeof(GetXsecResultStruct), cudaMemcpyDeviceToHost);
GetXsecResultStruct res = *(h_res);
if (res.y != -1) {
return res.y;
- }
+ }
else {
- G4double y = theInt.Interpolate(theManager.GetScheme(res.indexHigh), e,
+ G4double y = theInt.Interpolate(theManager.GetScheme(res.indexHigh), e,
res.pointLow.energy, res.pointHigh.energy,
res.pointLow.xSec, res.pointHigh.xSec);
if (nEntries == 1) {
@@ -569,7 +575,7 @@ void G4ParticleHPVector_CUDA::ThinOut(G4double precision) {
G4ParticleHPDataPoint *localTheData = (G4ParticleHPDataPoint*)malloc(nEntries * sizeof(G4ParticleHPDataPoint));
cudaMemcpy(localTheData, d_theData, nEntries * sizeof(G4ParticleHPDataPoint), cudaMemcpyDeviceToHost);
G4ParticleHPDataPoint * aBuff = new G4ParticleHPDataPoint[nPoints];
-
+
G4double x, x1, x2, y, y1, y2;
G4int count = 0, current = 2, start = 1;
@@ -628,7 +634,7 @@ __global__ void SampleGetResult_CUDA(G4ParticleHPDataPoint * theData, G4double *
G4double myRand;
G4double value;
G4double test;
-
+
G4int jcounter = 0;
G4int jcounter_max = 1024;
do {
@@ -639,17 +645,17 @@ __global__ void SampleGetResult_CUDA(G4ParticleHPDataPoint * theData, G4double *
}
myRand = rand_CUDA();
G4int ibin = SampleGetFirstIndex_CUDA(theIntegral, myRand, nEntries);
-
+
if (ibin < 0) {
printf("TKDB 080807 %f\n", myRand);
}
-
- // result
+
+ // result
myRand = rand_CUDA();
G4double x1, x2;
if (ibin == 0) {
- x1 = theData[ibin].energy;
- value = x1;
+ x1 = theData[ibin].energy;
+ value = x1;
break;
}
else {
@@ -658,20 +664,20 @@ __global__ void SampleGetResult_CUDA(G4ParticleHPDataPoint * theData, G4double *
x2 = theData[ibin].energy;
value = myRand * (x2 - x1) + x1;
-
+
// EMendoza - Always linear interpolation:
G4double y1 = theData[ibin-1].xSec;
G4double y2 = theData[ibin].xSec;
G4double mval = (y2-y1) / (x2-x1);
G4double bval = y1 - mval * x1;
- test = (mval * value + bval) / max(theData[ibin-1].xSec, theData[ibin].xSec);
+ test = (mval * value + bval) / max(theData[ibin-1].xSec, theData[ibin].xSec);
} while (rand_CUDA() > test);
*(result) = value;
}
-G4double G4ParticleHPVector_CUDA::Sample() {
+G4double G4ParticleHPVector_CUDA::Sample() {
G4double result;
-
+
int nBlocks = GetNumBlocks(nEntries);
SetAllNegativeXsecToZero_CUDA<<>> (d_theData, nEntries);
@@ -679,8 +685,8 @@ G4double G4ParticleHPVector_CUDA::Sample() {
cudaMemcpy(&result, &d_theData[0].energy, sizeof(G4double), cudaMemcpyHostToDevice);
}
else {
- if (d_theIntegral == 0) {
- IntegrateAndNormalise();
+ if (d_theIntegral == 0) {
+ IntegrateAndNormalise();
}
SampleGetResult_CUDA<<<1, 1>>> (d_theData, d_theIntegral, nEntries, d_singleDoubleResult);
cudaMemcpy(&result, d_singleDoubleResult, sizeof(G4double), cudaMemcpyDeviceToHost);
@@ -714,11 +720,11 @@ void G4ParticleHPVector_CUDA::Check(G4int i) {
int nBlocks = GetNumBlocks(nEntries);
CopyDataPointsToBuffer_CUDA<<>> (d_theData, d_newTheData, nEntries);
-
+
cudaFree(d_theData);
d_theData = d_newTheData;
}
-
+
if (i == nEntries) {
nEntries = i + 1;
}
@@ -731,5 +737,5 @@ G4bool G4ParticleHPVector_CUDA::IsBlocked(G4double aX) {
}
G4double G4ParticleHPVector_CUDA::GetUniformRand() {
- return (G4double)rand() / (G4double)RAND_MAX;
+ return (G4double)rand() / (G4double)RAND_MAX;
}