-
Notifications
You must be signed in to change notification settings - Fork 4
feat: add per-process and per-photon G4 timing instrumentation #280
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
base: main
Are you sure you want to change the base?
Changes from all commits
9aa9434
9bfd6fa
215f5ba
9535b89
bfdf3bb
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,114 @@ | ||
| # GPU vs G4 Optical Photon Benchmark | ||
|
|
||
| Measures the speedup of GPU (OptiX 7) optical photon propagation compared to | ||
| Geant4 CPU propagation on the APEX liquid argon scintillator tile detector. | ||
|
|
||
| ## Quick Start | ||
|
|
||
| ```bash | ||
| cd /workspaces/eic-opticks # or wherever the repo is | ||
| ./examples/benchmark_apex.sh # single event (~250k photons) | ||
| ./examples/benchmark_apex.sh 10 # 10 events (~2.5M photons) | ||
| ``` | ||
|
|
||
| Requires `apex.gdml` in the current directory and a working eic-opticks installation. | ||
|
|
||
| ## What It Measures | ||
|
|
||
| The benchmark runs GPURaytrace three times to isolate each timing component: | ||
|
|
||
| | Run | G4 photons | GPU | What it measures | | ||
| |-----|-----------|-----|------------------| | ||
| | 1. `--skip-gpu` | Yes | No | G4 electron tracking + G4 photon propagation | | ||
| | 2. `setStackPhotons false` | No | No | G4 electron tracking only (baseline) | | ||
| | 3. Normal | Yes | Yes | GPU photon propagation time | | ||
|
|
||
| **Speedup = (Run 1 - Run 2) / GPU time from Run 3** | ||
|
|
||
| - Run 1 uses `--skip-gpu` to prevent the GPU `simulate()` call, so the G4 wall | ||
| time reflects only CPU-side work. | ||
| - Run 2 uses G4's `setStackPhotons false` to prevent optical photon track creation | ||
| entirely — this is cheaper than killing photons in the SteppingAction because no | ||
| G4Track objects are allocated. | ||
| - GPU time is the wall-clock duration of `gx->simulate() + cudaDeviceSynchronize()`, | ||
| which is pure photon propagation on the GPU. | ||
|
|
||
| Runs 1 and 2 use single-threaded G4 (`/run/numberOfThreads 1`) to avoid | ||
| multithreading overhead inflating the per-photon CPU time. | ||
|
|
||
| ## Per-Process Timing | ||
|
|
||
| Run 1 also prints per-process step timing and per-photon statistics: | ||
|
|
||
| ``` | ||
| Geant4: StepTime Transportation: count= 114186428 avg= 0.77 us total= 87.840 s | ||
| Geant4: StepTime OpWLS: count= 42418 avg= 0.85 us total= 0.036 s | ||
| Geant4: StepTime OpRayleigh: count= 96947 avg= 0.98 us total= 0.095 s | ||
| Geant4: StepTime OpAbsorption: count= 8621 avg= 0.82 us total= 0.007 s | ||
|
|
||
| Geant4: PhotonTiming: 288738 photons, avg=306.0 us, median=2.1 us, ... | ||
| Geant4: PhotonPercentiles: p10=1.1 us, p50=2.1 us, p90=9.7 us, p99=7906.9 us | ||
| Geant4: PhotonSteps: 288738 photons, avg=397.0, median=2, min=1, max=10001 | ||
| Geant4: StepPercentiles: p10=1, p50=2, p90=10, p99=10001 | ||
| ``` | ||
|
|
||
| ### How step timing works | ||
|
|
||
| Each optical photon step is timed using `std::chrono::steady_clock` between | ||
| consecutive `UserSteppingAction` calls. The interval captures the full step cost: | ||
| geometry navigation (`G4Navigator::ComputeStep`), all process interaction length | ||
| proposals, and physics execution. Steps are classified by the post-step defining | ||
| process name. | ||
|
|
||
| ### How per-photon timing works | ||
|
|
||
| `PreUserTrackingAction` starts a timer when an optical photon track begins. | ||
| `PostUserTrackingAction` stops it and records the duration and step count. | ||
| Median and percentiles are computed from the collected per-photon times at end | ||
| of run. | ||
|
|
||
| Both approaches are thread-safe (thread-local timers, atomic accumulators). | ||
|
|
||
| ## Typical Results | ||
|
|
||
| ### APEX detector, 10 MeV electron, single event (~250k photons) | ||
|
|
||
| ``` | ||
| G4 photon-only: 88.4 s (306 us/photon avg, 2.1 us median) | ||
| GPU: 0.064 s (82 ns/photon) | ||
| Speedup: ~1,400x | ||
| ``` | ||
|
|
||
| ### APEX detector, 10 events (~2.5M photons) | ||
|
|
||
| ``` | ||
| G4 photon-only: ~880 s | ||
| GPU: 0.20 s (82 ns/photon) | ||
| Speedup: ~4,400x | ||
| ``` | ||
|
|
||
| GPU rate improves with batch size (better RTX occupancy). | ||
|
|
||
| ## Configuration | ||
|
|
||
| Environment variables: | ||
|
|
||
| | Variable | Default | Description | | ||
| |----------|---------|-------------| | ||
| | `OPTICKS_PROPAGATE_EPSILON` | 0.00001 | Ray tmin after boundary events (mm) | | ||
| | `OPTICKS_PROPAGATE_EPSILON0` | 0.0006 | Ray tmin after bulk events (mm) | | ||
| | `OPTICKS_MAX_BOUNCE` | 1000 | Maximum bounces per photon | | ||
| | `BENCHMARK_CONFIG` | det_debug | Config JSON name | | ||
|
|
||
| ## Implementation Details | ||
|
|
||
| The timing instrumentation is in `src/GPURaytrace.h`: | ||
|
|
||
| - **SteppingAction**: Atomic counters `fTimeTransport`, `fTimeOpWLS`, etc. | ||
| Thread-local `fLastStepTime` for inter-step interval measurement. | ||
| - **TrackingAction**: Atomic counters `fTotalPhotonTime`, `fPhotonCount`, etc. | ||
| Thread-local `fTrackStart` for per-photon lifetime measurement. | ||
| Mutex-protected vectors `fAllTimes`, `fAllSteps` for median/percentile computation. | ||
| - **RunAction**: `fSkipGPU` flag prevents `gx->simulate()`. `PrintTimingReport()` | ||
| outputs all collected timing data at end of run. | ||
| - **GPURaytrace.cpp**: `--skip-gpu` CLI flag sets `RunAction::fSkipGPU`. | ||
|
Member
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. This new script looks like a benchmark measurement. Let's move it to benchmarks/ |
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,128 @@ | ||
| #!/bin/bash | ||
| # photontimingandsteps.sh — Measure GPU vs G4 optical photon speedup and timing | ||
| # | ||
| # Methodology (3 runs): | ||
| # | ||
| # Run 1 (--skip-gpu): G4 tracks electron + propagates optical photons | ||
| # GPU is skipped. Genstep collection is skipped. | ||
| # Measures: G4 photon-only CPU time, per-process | ||
| # step timing, per-photon timing with percentiles. | ||
| # | ||
| # Run 2 (setStackPhotons off): G4 tracks electron only, no photon tracks created. | ||
| # GPU is skipped. | ||
| # Measures: EM baseline time (electron tracking only). | ||
| # | ||
| # Run 3 (setStackPhotons off, G4 tracks electron and collects gensteps but does | ||
| # GPU active): NOT create photon tracks. GPU propagates photons. | ||
| # Measures: GPU-accelerated total wall time. | ||
| # | ||
| # Speedup calculations: | ||
| # G4 photon-only CPU time = Run 1 CPU - Run 2 CPU | ||
| # Photon-only speedup = G4 photon-only / GPU sim time | ||
| # Wall time speedup = Run 1 wall / Run 3 wall | ||
| # | ||
| # Usage: | ||
| # ./examples/photontiming_geant4/photontimingandsteps.sh [events] # default: 1 | ||
|
|
||
| set -e | ||
|
|
||
| GDML="apex.gdml" | ||
| EPS="${OPTICKS_PROPAGATE_EPSILON:-0.00001}" | ||
| EPS0="${OPTICKS_PROPAGATE_EPSILON0:-0.0006}" | ||
| CONFIG="${BENCHMARK_CONFIG:-det_debug}" | ||
| NEVENTS="${1:-1}" | ||
|
|
||
| if [ ! -f "$GDML" ]; then | ||
| echo "ERROR: $GDML not found. Run from the eic-opticks root directory." | ||
| exit 1 | ||
| fi | ||
|
|
||
| echo "=== Optical Photon Timing & Speedup Benchmark ===" | ||
| echo "GDML=$GDML eps=$EPS eps0=$EPS0 events=$NEVENTS config=$CONFIG" | ||
|
|
||
| # Macro: normal G4 (photons propagated) | ||
| MAC_NORMAL=$(mktemp /tmp/bench_normal_XXXXXX.mac) | ||
| cat > "$MAC_NORMAL" << EOF | ||
| /run/verbose 1 | ||
| /run/numberOfThreads 1 | ||
| /run/initialize | ||
| /run/beamOn $NEVENTS | ||
| EOF | ||
|
|
||
| # Macro: no photon tracks (gensteps still collected by SteppingAction) | ||
| MAC_NOPHOTON=$(mktemp /tmp/bench_nophoton_XXXXXX.mac) | ||
| cat > "$MAC_NOPHOTON" << EOF | ||
| /run/verbose 1 | ||
| /run/numberOfThreads 1 | ||
| /run/initialize | ||
| /process/optical/scintillation/setStackPhotons false | ||
| /process/optical/cerenkov/setStackPhotons false | ||
| /run/beamOn $NEVENTS | ||
| EOF | ||
|
|
||
| export OPTICKS_MAX_BOUNCE=1000 | ||
| export OPTICKS_PROPAGATE_EPSILON=$EPS | ||
| export OPTICKS_PROPAGATE_EPSILON0=$EPS0 | ||
|
|
||
| # --------------------------------------------------------------- | ||
| # Run 1: G4 photon propagation (no GPU, no genstep collection) | ||
| # --------------------------------------------------------------- | ||
| echo "" | ||
| echo "--- Run 1: G4 photon propagation (--skip-gpu, single-threaded) ---" | ||
| LOG1=$(mktemp /tmp/bench1_XXXXXX.txt) | ||
| GPURaytrace -g "$GDML" -m "$MAC_NORMAL" -c "$CONFIG" -s 42 --skip-gpu &> "$LOG1" | ||
|
|
||
| G4_LINE=$(grep "^ User=" "$LOG1" | tail -1) | ||
| G4_PHOTON_CPU=$(echo "$G4_LINE" | grep -oP 'User=\K[0-9.]+') | ||
| G4_PHOTON_WALL=$(echo "$G4_LINE" | grep -oP 'Real=\K[0-9.]+') | ||
| G4_HITS=$(grep "Geant4: NumHits:" "$LOG1" | awk '{print $NF}') | ||
|
|
||
| echo " G4 CPU time: ${G4_PHOTON_CPU}s" | ||
| echo " G4 wall time: ${G4_PHOTON_WALL}s" | ||
| echo " G4 hits: $G4_HITS" | ||
|
|
||
| # Print per-process and per-photon timing | ||
| grep "Geant4: StepTime\|Geant4: PhotonTiming\|Geant4: PhotonPercentiles\|Geant4: PhotonSteps\|Geant4: StepPercentiles\|Geant4: PhotonTimeHist\|Geant4: OpticalSteps" "$LOG1" | sed 's/^/ /' | ||
|
|
||
| # --------------------------------------------------------------- | ||
| # Run 2: EM baseline (no photon tracks, no GPU) | ||
| # --------------------------------------------------------------- | ||
| echo "" | ||
| echo "--- Run 2: EM baseline (setStackPhotons false, --skip-gpu) ---" | ||
| LOG2=$(mktemp /tmp/bench2_XXXXXX.txt) | ||
| GPURaytrace -g "$GDML" -m "$MAC_NOPHOTON" -c "$CONFIG" -s 42 --skip-gpu &> "$LOG2" | ||
|
|
||
| G4_BASE_LINE=$(grep "^ User=" "$LOG2" | tail -1) | ||
| G4_BASE_CPU=$(echo "$G4_BASE_LINE" | grep -oP 'User=\K[0-9.]+') | ||
| echo " EM baseline CPU: ${G4_BASE_CPU}s" | ||
|
|
||
| # --------------------------------------------------------------- | ||
| # Run 3: GPU-accelerated (no G4 photon tracks, GPU propagates) | ||
| # --------------------------------------------------------------- | ||
| echo "" | ||
| echo "--- Run 3: GPU-accelerated (setStackPhotons false, GPU active) ---" | ||
| LOG3=$(mktemp /tmp/bench3_XXXXXX.txt) | ||
| GPURaytrace -g "$GDML" -m "$MAC_NOPHOTON" -c "$CONFIG" -s 42 &> "$LOG3" | ||
|
|
||
| GPU_TIME=$(grep "Simulation time:" "$LOG3" | awk '{print $3}') | ||
| GPU_HITS=$(grep "Opticks: NumHits:" "$LOG3" | awk '{print $NF}') | ||
| GPU_LINE=$(grep "^ User=" "$LOG3" | tail -1) | ||
| GPU_WALL=$(echo "$GPU_LINE" | grep -oP 'Real=\K[0-9.]+') | ||
| NPHOTONS=$(grep "NumCollected:" "$LOG3" | tail -1 | awk '{print $NF}') | ||
|
|
||
| echo " GPU sim time: ${GPU_TIME}s" | ||
| echo " GPU total wall: ${GPU_WALL}s" | ||
| echo " GPU hits: $GPU_HITS" | ||
| echo " Photons: $NPHOTONS" | ||
|
|
||
| # --------------------------------------------------------------- | ||
| # Results | ||
| # --------------------------------------------------------------- | ||
| echo "" | ||
| echo "=== Results ===" | ||
| python3 optiphy/ana/photon_timing_summary.py \ | ||
| "${GPU_TIME:-0}" "${G4_PHOTON_CPU:-0}" "${G4_BASE_CPU:-0}" \ | ||
| "${G4_PHOTON_WALL:-0}" "${GPU_WALL:-0}" \ | ||
| "${NPHOTONS:-0}" "${GPU_HITS:-0}" "${G4_HITS:-0}" | ||
|
|
||
| rm -f "$LOG1" "$LOG2" "$LOG3" "$MAC_NORMAL" "$MAC_NOPHOTON" |
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,52 @@ | ||
| """ | ||
| photon_timing_summary.py | ||
| ======================== | ||
| Format and print the photon-only and wall-time speedup summary used by | ||
| examples/photontiming_geant4/photontimingandsteps.sh after the three | ||
| benchmark runs (G4 photon-only, EM baseline, GPU-accelerated) complete. | ||
|
|
||
| Usage: | ||
| python3 photon_timing_summary.py <gpu_sim> <g4_cpu> <g4_base> <g4_wall> \ | ||
| <gpu_wall> <nphotons> <gpu_hits> <g4_hits> | ||
| """ | ||
|
|
||
| import sys | ||
|
|
||
|
|
||
| def main(argv): | ||
| gpu_sim = float(argv[1]) | ||
| g4_cpu = float(argv[2]) | ||
| g4_base = float(argv[3]) | ||
| g4_wall = float(argv[4]) | ||
| gpu_wall = float(argv[5]) | ||
| nphotons = int(argv[6]) | ||
| gpu_hits = int(argv[7]) | ||
| g4_hits = int(argv[8]) | ||
|
|
||
| g4_photon_cpu = g4_cpu - g4_base | ||
| hit_diff = (gpu_hits - g4_hits) / g4_hits * 100 if g4_hits > 0 else 0 | ||
|
|
||
| print() | ||
| print(f'Photons: {nphotons:>12,}') | ||
| print() | ||
| print(f'--- Photon-only speedup ---') | ||
| print(f'G4 photon CPU time: {g4_photon_cpu:>12.2f} s') | ||
| print(f'GPU sim time: {gpu_sim:>12.4f} s') | ||
| if gpu_sim > 0 and g4_photon_cpu > 0: | ||
| print(f'Photon speedup: {g4_photon_cpu/gpu_sim:>12.0f}x') | ||
| print(f'GPU time/photon: {gpu_sim/nphotons*1e9:>12.1f} ns') | ||
| print(f'G4 time/photon (avg): {g4_photon_cpu/nphotons*1e6:>12.1f} us') | ||
| print() | ||
| print(f'--- Wall time speedup ---') | ||
| print(f'G4 wall (EM + photons): {g4_wall:>12.2f} s') | ||
| print(f'GPU wall (EM + GPU): {gpu_wall:>12.4f} s') | ||
| if gpu_wall > 0 and g4_wall > 0: | ||
| print(f'Wall speedup: {g4_wall/gpu_wall:>12.0f}x') | ||
| print() | ||
| print(f'GPU hits: {gpu_hits:>12}') | ||
| print(f'G4 hits: {g4_hits:>12}') | ||
| print(f'Hit diff: {hit_diff:>+11.1f}%') | ||
|
|
||
|
|
||
| if __name__ == "__main__": | ||
| main(sys.argv) |
| Original file line number | Diff line number | Diff line change | ||||||||||||||||||||||||
|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|
|
|
@@ -75,6 +75,10 @@ int main(int argc, char **argv) | |||||||||||||||||||||||||
|
|
||||||||||||||||||||||||||
| program.add_argument("-s", "--seed").help("fixed random seed (default: time-based)").scan<'i', long>(); | ||||||||||||||||||||||||||
|
|
||||||||||||||||||||||||||
| program.add_argument("--skip-gpu") | ||||||||||||||||||||||||||
| .help("skip GPU photon propagation (for measuring G4-only photon time)") | ||||||||||||||||||||||||||
| .flag(); | ||||||||||||||||||||||||||
|
Comment on lines
+78
to
+80
Contributor
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. clang-format suggestion
Suggested change
Comment on lines
+78
to
+80
Contributor
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. clang-format suggestion
Suggested change
Comment on lines
+78
to
+80
Contributor
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. clang-format suggestion
Suggested change
plexoos marked this conversation as resolved.
|
||||||||||||||||||||||||||
|
|
||||||||||||||||||||||||||
| try | ||||||||||||||||||||||||||
| { | ||||||||||||||||||||||||||
| program.parse_args(argc, argv); | ||||||||||||||||||||||||||
|
|
@@ -108,6 +112,13 @@ int main(int argc, char **argv) | |||||||||||||||||||||||||
|
|
||||||||||||||||||||||||||
| G4App *g4app = new G4App(gdml_file); | ||||||||||||||||||||||||||
|
|
||||||||||||||||||||||||||
| if (program.get<bool>("--skip-gpu")) | ||||||||||||||||||||||||||
| { | ||||||||||||||||||||||||||
| g4app->run_act_->fSkipGPU = true; | ||||||||||||||||||||||||||
| g4app->stepping_->fSkipGenstep = true; | ||||||||||||||||||||||||||
| G4cout << "Skip-GPU mode: GPU propagation and genstep collection skipped" << G4endl; | ||||||||||||||||||||||||||
| } | ||||||||||||||||||||||||||
|
|
||||||||||||||||||||||||||
| ActionInitialization *actionInit = new ActionInitialization(g4app); | ||||||||||||||||||||||||||
| run_mgr->SetUserInitialization(actionInit); | ||||||||||||||||||||||||||
| run_mgr->SetUserInitialization(g4app->det_cons_); | ||||||||||||||||||||||||||
|
|
||||||||||||||||||||||||||
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Is this benchmark really for OptiX 7 only?