Tweet
Tweet

Slide 1

Slide 1 text

Mitglied der Helmholtz-Gemeinschaft 1 GPU in High Energy Physics Conference, Pisa 11 September 2014, Andreas Herten GPU-based Online Tracking for the PANDA Experiment

Slide 2

Slide 2 text

Mitglied der Helmholtz-Gemeinschaft Outline • PANDA – Experiment – Online Event Filter • Algorithms – Hough Transform – Riemann Track Finder – Triplet Finder 2

Slide 3

Slide 3 text

Mitglied der Helmholtz-Gemeinschaft FAIR • Facility for Antiproton and Ion Research – New accelerator complex (Darmstadt, Germany) – Next to GSI laboratory – Construction in progress, ending 2018 – Four pillars of research: 3 APPA NUSTAR CBM PANDA Atom & plasma physics Nuclear structure, astro physics Hadron physics Hadron physics

Slide 4

Slide 4 text

Mitglied der Helmholtz-Gemeinschaft FAIR • Facility for Antiproton and Ion Research – New accelerator complex (Darmstadt, Germany) – Next to GSI laboratory – Construction in progress, ending 2018 – Four pillars of research: 3 APPA NUSTAR CBM PANDA Atom & plasma physics Nuclear structure, astro physics Hadron physics Hadron physics © Google Maps

Slide 5

Slide 5 text

Mitglied der Helmholtz-Gemeinschaft FAIR • Facility for Antiproton and Ion Research – New accelerator complex (Darmstadt, Germany) – Next to GSI laboratory – Construction in progress, ending 2018 – Four pillars of research: 3 APPA NUSTAR CBM PANDA Atom & plasma physics Nuclear structure, astro physics Hadron physics Hadron physics

Slide 6

Slide 6 text

Mitglied der Helmholtz-Gemeinschaft FAIR • Facility for Antiproton and Ion Research – New accelerator complex (Darmstadt, Germany) – Next to GSI laboratory – Construction in progress, ending 2018 – Four pillars of research: 3 APPA NUSTAR CBM PANDA Atom & plasma physics Nuclear structure, astro physics Hadron physics Hadron physics fair-center.eu

Slide 7

Slide 7 text

Mitglied der Helmholtz-Gemeinschaft Mitglied der Helmholtz-Gemeinschaft FAIR Accelerator Complex 4 existing facilities new facilities experiments

Slide 8

Slide 8 text

Mitglied der Helmholtz-Gemeinschaft Mitglied der Helmholtz-Gemeinschaft FAIR Accelerator Complex 4 existing facilities new facilities experiments Beam momentum: p = 1.5 - 15 GeV/c High resolution: L = 1031 cm-2s-1 Δp/p < 4·10-5 High luminosity: L = 2·1032 cm-2s-1 Δp/p < 10-4 High Energy Storage Ring

Slide 9

Slide 9 text

Mitglied der Helmholtz-Gemeinschaft PANDA — The Experiment 5 13 m p p

Slide 10

Slide 10 text

Mitglied der Helmholtz-Gemeinschaft PANDA — The Experiment 5 13 m p p Magnet STT MVD

Slide 11

Slide 11 text

Mitglied der Helmholtz-Gemeinschaft PANDA — MVD & STT Micro Vertex Detector • Silicon-based pixel + strip detector • 10 000 000 + 200 000 channels • Vertex resolution: < 100 μm 6 Straw Tube Tracker • 4636 small drift tubes (ø 1 cm) • Drift times: < 250 ns • 26 layers, 8 skewed • Material budget: 1.2 % radiation length

Slide 12

Slide 12 text

Mitglied der Helmholtz-Gemeinschaft PANDA — Physics • Meson spectroscopy – Light mesons – Charmonium – Exotic states • Glueballs • Hybrids • Molecules/multiquarks – Open charm • Baryon production • Nucleon structure, e.m. processes • Charm in nuclei • Strangeness physics 7 0 2 4 6 8 12 15 10 p Momentum / GeV/c Mass / GeV/c2 1 2 3 4 5 6 ΛΛ ΣΣ ΞΞ Λ c Λ c Σ c Σ c Ξ c Ξ c Ω c Ω c ΩΩ DD D s D s ggg,gg light qq π,ρ,ω,f 2 ,K,K* cc J/ψ, η c , χ cJ qqqq ccqq nng,ssg ccg nng,ssg ccg ggg

Slide 13

Slide 13 text

Mitglied der Helmholtz-Gemeinschaft PANDA — Physics • Meson spectroscopy – Light mesons – Charmonium – Exotic states • Glueballs • Hybrids • Molecules/multiquarks – Open charm • Baryon production • Nucleon structure, e.m. processes • Charm in nuclei • Strangeness physics 7 → Broad physics program 0 2 4 6 8 12 15 10 p Momentum / GeV/c Mass / GeV/c2 1 2 3 4 5 6 ΛΛ ΣΣ ΞΞ Λ c Λ c Σ c Σ c Ξ c Ξ c Ω c Ω c ΩΩ DD D s D s ggg,gg light qq π,ρ,ω,f 2 ,K,K* cc J/ψ, η c , χ cJ qqqq ccqq nng,ssg ccg nng,ssg ccg ggg

Slide 14

Slide 14 text

Mitglied der Helmholtz-Gemeinschaft PANDA — Event Reconstruction • Continuous read out – Background & signal similar – Novel feature – No hardware trigger based on few sub-detectors, but online event reconstruction using full detector information 8 (Reject background events, save interesting events) Reduction Amount: Time: ~1/1000 50 ns/evt Storage space for offline analysis 3 PB/y Event: Raw data: 2 × 107/s 200 GB/s Rate

Slide 15

Slide 15 text

Mitglied der Helmholtz-Gemeinschaft PANDA — Read Out Scheme 9

Slide 16

Slide 16 text

Mitglied der Helmholtz-Gemeinschaft PANDA — Read Out Scheme • Requirements to Online Tracking • Fast • Sophisticated algorithms possible; reprogrammable • Parallelity beyond single devices • Fast • Limited precision ok 9

Slide 17

Slide 17 text

Mitglied der Helmholtz-Gemeinschaft PANDA — Read Out Scheme • Requirements to Online Tracking • Fast • Sophisticated algorithms possible; reprogrammable • Parallelity beyond single devices • Fast • Limited precision ok 9 GPUs

Slide 18

Slide 18 text

Mitglied der Helmholtz-Gemeinschaft ALGORITHMS #1 10

Slide 19

Slide 19 text

Mitglied der Helmholtz-Gemeinschaft ALGORITHMS #1 11 Hough Transform Riemann Track Finder Triplet Finder

Slide 20

Slide 20 text

Mitglied der Helmholtz-Gemeinschaft Algorithm: Hough Transform • Idea: Transform (x,y)i → (α,r)ij , find lines via (α,r) space • Solve rij line equation for – Lots of hits (x,y,ρ)i and – Many αj ∈ [0°,360°) each • Fill histogram • Extract track parameters 12 rij = cos ↵j · xi + sin ↵j · yi + ⇢i i: ~100 hits/event (STT) j: every 0.2° rij : 180 000 x y x y Mitglied der Helmholtz-Gemeinschaft Hough Transform — Princip → Bin giv r α

Slide 21

Slide 21 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Visualization Lines 14 • Create lines going through hit point (x,y)i – Line parameterized by rij = cos(αj ) ⋅ xi + sin(αj ) ⋅ yi • Fill line parameters (α,r)ij into histogram – Rasterize for many αj ∈ [0°,180°) α = 0°, 10°, 20°, …

Slide 22

Slide 22 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Visualization Lines 14 • Create lines going through hit point (x,y)i – Line parameterized by rij = cos(αj ) ⋅ xi + sin(αj ) ⋅ yi • Fill line parameters (α,r)ij into histogram – Rasterize for many αj ∈ [0°,180°) α = 0°, 10°, 20°, …

Slide 23

Slide 23 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Visualization Lines 14 • Create lines going through hit point (x,y)i – Line parameterized by rij = cos(αj ) ⋅ xi + sin(αj ) ⋅ yi • Fill line parameters (α,r)ij into histogram – Rasterize for many αj ∈ [0°,180°) α = 0°, 10°, 20°, …

Slide 24

Slide 24 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Visualization Lines 14 • Create lines going through hit point (x,y)i – Line parameterized by rij = cos(αj ) ⋅ xi + sin(αj ) ⋅ yi • Fill line parameters (α,r)ij into histogram – Rasterize for many αj ∈ [0°,180°) α = 0°, 10°, 20°, …

Slide 25

Slide 25 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Visualization Points 16 • Create lines going through hit point (x,y)i • Repeat for every hit point i α = 0°, 10°, 20°, …

Slide 26

Slide 26 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Visualization Points 16 • Create lines going through hit point (x,y)i • Repeat for every hit point i α = 0°, 10°, 20°, …

Slide 27

Slide 27 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Visualization Points 16 • Create lines going through hit point (x,y)i • Repeat for every hit point i α = 0°, 10°, 20°, …

Slide 28

Slide 28 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Visualization Points 16 • Create lines going through hit point (x,y)i • Repeat for every hit point i α = 0°, 10°, 20°, …

Slide 29

Slide 29 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Granularity 18 • Choice of α granularity determines resolution α = 0°, 2°, 4°, … α = 0°, 2°, 4°, … i: ~100 hits/event (STT) j: every 0.2° rij : 180 000

Slide 30

Slide 30 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Granularity 18 • Choice of α granularity determines resolution α = 0°, 2°, 4°, … α = 0°, 2°, 4°, … i: ~100 hits/event (STT) j: every 0.2° rij : 180 000

Slide 31

Slide 31 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Granularity 18 • Choice of α granularity determines resolution α = 0°, 2°, 4°, … α = 0°, 2°, 4°, … i: ~100 hits/event (STT) j: every 0.2° rij : 180 000

Slide 32

Slide 32 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Granularity 18 • Choice of α granularity determines resolution α = 0°, 2°, 4°, … α = 0°, 2°, 4°, … i: ~100 hits/event (STT) j: every 0.2° rij : 180 000

Slide 33

Slide 33 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Granularity 18 • Choice of α granularity determines resolution α = 0°, 2°, 4°, … α = 0°, 2°, 4°, … i: ~100 hits/event (STT) j: every 0.2° rij : 180 000

Slide 34

Slide 34 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Granularity 18 • Choice of α granularity determines resolution α = 0°, 2°, 4°, … α = 0°, 2°, 4°, … i: ~100 hits/event (STT) j: every 0.2° rij : 180 000

Slide 35

Slide 35 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Granularity 18 • Choice of α granularity determines resolution α = 0°, 2°, 4°, … α = 0°, 2°, 4°, … i: ~100 hits/event (STT) j: every 0.2° rij : 180 000

Slide 36

Slide 36 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Granularity 18 • Choice of α granularity determines resolution α = 0°, 2°, 4°, … α = 0°, 2°, 4°, … i: ~100 hits/event (STT) j: every 0.2° rij : 180 000

Slide 37

Slide 37 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Granularity 18 • Choice of α granularity determines resolution α = 0°, 2°, 4°, … α = 0°, 2°, 4°, … i: ~100 hits/event (STT) j: every 0.2° rij : 180 000

Slide 38

Slide 38 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Granularity 18 • Choice of α granularity determines resolution α = 0°, 2°, 4°, … α = 0°, 2°, 4°, … i: ~100 hits/event (STT) j: every 0.2° rij : 180 000

Slide 39

Slide 39 text

° Angle / 0 20 40 60 80 100 120 140 160 180 Hough transformed -0.4 -0.3 -0.2 -0.1 0 0.1 0.2 0.3 0.4 0.5 0.6 0 Entries 2.2356e+08 Mean x 90 Mean y 0.02905 RMS x 51.96 RMS y 0.1063 0 5 10 15 20 25 0 Entries 2.2356e+08 Mean x 90 Mean y 0.02905 RMS x 51.96 RMS y 0.1063 1800 x 1800 Grid PANDA STT+MVD Mitglied der Helmholtz-Gemeinschaft 19 68 (x,y) points r α Algorithm: Hough Transform

Slide 40

Slide 40 text

° Angle / 0 20 40 60 80 100 120 140 160 180 Hough transformed -0.4 -0.3 -0.2 -0.1 0 0.1 0.2 0.3 0.4 0.5 0.6 0 Entries 2.2356e+08 Mean x 90 Mean y 0.02905 RMS x 51.96 RMS y 0.1063 0 5 10 15 20 25 0 Entries 2.2356e+08 Mean x 90 Mean y 0.02905 RMS x 51.96 RMS y 0.1063 1800 x 1800 Grid PANDA STT+MVD Mitglied der Helmholtz-Gemeinschaft 19 68 (x,y) points r α Algorithm: Hough Transform

Slide 41

Slide 41 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Remarks Two Implementations • Thrust (CUDA‘s STL) – Performance: 3 ms/evt – Reduce to set of standard routines • Fast (uses Thrust‘s optimized algorithms) • Inflexible (hard to customize) • Not yet at performance maximum • Plain CUDA – Performance: 0.5 ms/evt – Build completely for this task • Fitting for PANDA; customizable • A bit more complicated at parts • 20 Peakfinding challenging

Slide 42

Slide 42 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Remarks 21 ° / α 0 20 40 60 80 100 120 140 160 180 r -30 -20 -10 0 10 20 30 40 HoughHist Entries 9000 Mean x 89.33 Mean y 6.66 RMS x 51.8 RMS y 19.2 0 2 4 6 8 10 12 14 16 18 HoughHist Entries 9000 Mean x 89.33 Mean y 6.66 RMS x 51.8 RMS y 19.2 HT histogram Hill Climber Peakfinding challenging

Slide 43

Slide 43 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Remarks 21 ° / α 0 20 40 60 80 100 120 140 160 180 r -30 -20 -10 0 10 20 30 40 houghIt0 Entries 9000 Mean x 89.33 Mean y 6.66 RMS x 51.8 RMS y 19.2 0 2 4 6 8 10 12 14 16 18 houghIt0 Entries 9000 Mean x 89.33 Mean y 6.66 RMS x 51.8 RMS y 19.2 HT histogram ° / α 0 20 40 60 80 100 120 140 160 180 r -30 -20 -10 0 10 20 30 40 houghIt1 Entries 5580 Mean x 89.6 Mean y 9.719 RMS x 51.78 RMS y 18.09 0 2 4 6 8 10 12 14 16 houghIt1 Entries 5580 Mean x 89.6 Mean y 9.719 RMS x 51.78 RMS y 18.09 HT histogram ° / α 0 20 40 60 80 100 120 140 160 180 r -30 -20 -10 0 10 20 30 houghIt2 Entries 2700 Mean x 89.13 Mean y 13.79 RMS x 51.77 RMS y 14.04 0 2 4 6 8 10 12 houghIt2 Entries 2700 Mean x 89.13 Mean y 13.79 RMS x 51.77 RMS y 14.04 HT histogram -40 -30 -20 -10 0 10 20 30 40 0 5 10 15 20 25 30 Iterative Maximum Deleter Peakfinding challenging

Slide 44

Slide 44 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Remarks 21 ° / α 0 20 40 60 80 100 120 140 160 180 r -30 -20 -10 0 10 20 30 40 houghIt0 Entries 9000 Mean x 89.33 Mean y 6.66 RMS x 51.8 RMS y 19.2 0 2 4 6 8 10 12 14 16 18 houghIt0 Entries 9000 Mean x 89.33 Mean y 6.66 RMS x 51.8 RMS y 19.2 HT histogram ° / α 0 20 40 60 80 100 120 140 160 180 r -30 -20 -10 0 10 20 30 40 houghIt1 Entries 5580 Mean x 89.6 Mean y 9.719 RMS x 51.78 RMS y 18.09 0 2 4 6 8 10 12 14 16 houghIt1 Entries 5580 Mean x 89.6 Mean y 9.719 RMS x 51.78 RMS y 18.09 HT histogram ° / α 0 20 40 60 80 100 120 140 160 180 r -30 -20 -10 0 10 20 30 houghIt2 Entries 2700 Mean x 89.13 Mean y 13.79 RMS x 51.77 RMS y 14.04 0 2 4 6 8 10 12 houghIt2 Entries 2700 Mean x 89.13 Mean y 13.79 RMS x 51.77 RMS y 14.04 HT histogram -40 -30 -20 -10 0 10 20 30 40 0 5 10 15 20 25 30 Iterative Maximum Deleter Peakfinding challenging current research

Slide 45

Slide 45 text

Mitglied der Helmholtz-Gemeinschaft 22 ALGORITHMS #2 Hough Transform Riemann Track Finder Triplet Finder

Slide 46

Slide 46 text

Mitglied der Helmholtz-Gemeinschaft 23 Riemann Track Finder — Method • Idea: Don‘t fit lines (in 2D), fit planes (in 3D)! • Create seeds – All possible three hit combinations • Grow seeds to tracks Continuously test next hit if it fits – Use mapping to Riemann paraboloid (+ s-z fit, det. layer) x x x x y z‘ x x x y x x x x y x More on: Seeds; Growing 1 2

Slide 47

Slide 47 text

Mitglied der Helmholtz-Gemeinschaft 24 Riemann Track Finder — GPU Adaptations CPU GPU 3 loops to generate seeds serially for (int i = 0; i < hitsInLayerOne.size(); i++) { for (int j = 0; j < hitsInLayerTwo.size(); j++) { for (int k = 0; k < hitsInLayerThree.size(); k++) { /* Triplet Generation */ } } } Needed: Mapping of inherent GPU indexing variable to triplet index int ijk = threadIdx.x + blockIdx.x * blockDim.x; nLayerx = 1 2 ⇣p 8x + 1 1 ⌘ pos ( nLayerx ) = 3 pp 3 p 243x2 1 + 27x 32 / 3 + 1 3 p 3 3 pp 3 p 243x2 1 + 27x 1 1 2 Port of CPU code; parallelism on seed base Only easy computations; e.g. 3x3 matrices

Slide 48

Slide 48 text

Mitglied der Helmholtz-Gemeinschaft 24 Riemann Track Finder — GPU Adaptations CPU GPU → 100 × faster than CPU version: ~0.6 ms/event 3 loops to generate seeds serially for (int i = 0; i < hitsInLayerOne.size(); i++) { for (int j = 0; j < hitsInLayerTwo.size(); j++) { for (int k = 0; k < hitsInLayerThree.size(); k++) { /* Triplet Generation */ } } } Needed: Mapping of inherent GPU indexing variable to triplet index int ijk = threadIdx.x + blockIdx.x * blockDim.x; nLayerx = 1 2 ⇣p 8x + 1 1 ⌘ pos ( nLayerx ) = 3 pp 3 p 243x2 1 + 27x 32 / 3 + 1 3 p 3 3 pp 3 p 243x2 1 + 27x 1 1 2 Port of CPU code; parallelism on seed base Only easy computations; e.g. 3x3 matrices

Slide 49

Slide 49 text

Mitglied der Helmholtz-Gemeinschaft 25 ALGORITHMS #3 Hough Transform Riemann Track Finder Triplet Finder

Slide 50

Slide 50 text

Mitglied der Helmholtz-Gemeinschaft 26 Triplet Finder • Algorithm specifically designed for the PANDA Straw Tube Tracker (STT) • Ported to GPU by Andrew Adinetz – CUDA, Dynamic Parallelism, Thrust – Quality of tracks comparable to CPU http://www.fz-juelich.de/ias/jsc/ Original algorithm by Marius Mertens et al 1.5 m

Slide 51

Slide 51 text

Mitglied der Helmholtz-Gemeinschaft 27 Triplet Finder • Idea: Use only subset of detector as seed – Don‘t use STT isochrones (drift times) – Calculate circle from 3 points (no fit) • Features – Fast & robust algorithm, no event time needed – Many tuning possibilities More

Slide 52

Slide 52 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Display 29 Triplet Isochrone early Isochrone early & skewed Isochrone close Isochrone late MVD hit Track timed out Track current

Slide 53

Slide 53 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Display 29 Triplet Isochrone early Isochrone early & skewed Isochrone close Isochrone late MVD hit Track timed out Track current

Slide 54

Slide 54 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Display 29 Triplet Isochrone early Isochrone early & skewed Isochrone close Isochrone late MVD hit Track timed out Track current

Slide 55

Slide 55 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Display 29 Triplet Isochrone early Isochrone early & skewed Isochrone close Isochrone late MVD hit Track timed out Track current

Slide 56

Slide 56 text

Mitglied der Helmholtz-Gemeinschaft 30 Triplet Finder — Times

Slide 57

Slide 57 text

Mitglied der Helmholtz-Gemeinschaft 30 Triplet Finder — Times

Slide 58

Slide 58 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Optimizations • Bunching Wrapper – Hits from one event have similar timestamp – Combine hits to sets (bunches) which occupy GPU best 31

Slide 59

Slide 59 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Optimizations • Bunching Wrapper – Hits from one event have similar timestamp – Combine hits to sets (bunches) which occupy GPU best 31 Hit

Slide 60

Slide 60 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Optimizations • Bunching Wrapper – Hits from one event have similar timestamp – Combine hits to sets (bunches) which occupy GPU best 31 Hit Event

Slide 61

Slide 61 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Optimizations • Bunching Wrapper – Hits from one event have similar timestamp – Combine hits to sets (bunches) which occupy GPU best 31 Hit Event

Slide 62

Slide 62 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Optimizations • Bunching Wrapper – Hits from one event have similar timestamp – Combine hits to sets (bunches) which occupy GPU best 31 Hit Event Bunch

Slide 63

Slide 63 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Optimizations • Bunching Wrapper – Hits from one event have similar timestamp – Combine hits to sets (bunches) which occupy GPU best 31 Hit Event Bunch !(N2) → !(N)

Slide 64

Slide 64 text

Mitglied der Helmholtz-Gemeinschaft 32 Triplet Finder — Bunching Performance

Slide 65

Slide 65 text

Dynamic Parallelism Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Optimizations • Compare kernel launch strategies 33 1 thread/bunch Calling kernel 1 thread/bunch Calling kernel Triplet Finder 1 thread/bunch Calling kernel 1 block/bunch Joined kernel 1 block/bunch Joined kernel 1 block/bunch Joined kernel TF Stage #1 TF Stage #2 TF Stage #3 TF Stage #4 1 stream/bunch Combining stream 1 stream/bunch Combining stream 1 stream/bunch Calling stream Joined Kernel Host Streams Triplet Finder Triplet Finder CPU GPU TF Stage #1 TF Stage #2 TF Stage #3 TF Stage #4 TF Stage #1 TF Stage #2 TF Stage #3 TF Stage #4

Slide 66

Slide 66 text

Mitglied der Helmholtz-Gemeinschaft 34 Triplet Finder — Kernel Launches Explanation

Slide 67

Slide 67 text

Mitglied der Helmholtz-Gemeinschaft 35 Triplet Finder — Clock Speed / GPU K40 3004 MHz, 745 MHz / 875 MHz K20X 2600 MHz, 732 MHz / 784 MHz Memory Clock Core Clock GPU Boost

Slide 68

Slide 68 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Summary • Best performance: 20 µs/event – 20⋅10-6 s/event * 2⋅107 event/s 㱺 400 GPUs2014 – PANDA2019: Multi GPU system – !(100) GPUs • Optimizations possible & needed – ε needs to be improved – Speed, €: • More float less double-cards a la K10 • Consumer-grade cards a la GTX 36

Slide 69

Slide 69 text

Mitglied der Helmholtz-Gemeinschaft Summary • PANDA researches in using GPUs as part of online event reconstruction scheme • Algorithms in active evaluation and optimization – Triplet Finder performance-optimized • Data transfer to GPU in research: FairMQ → Poster by Ludovico Bianchi 37

Slide 70

Slide 70 text

Thank you! Andreas Herten [email protected] Mitglied der Helmholtz-Gemeinschaft Summary • PANDA researches in using GPUs as part of online event reconstruction scheme • Algorithms in active evaluation and optimization – Triplet Finder performance-optimized • Data transfer to GPU in research: FairMQ → Poster by Ludovico Bianchi 37

Slide 71

Slide 71 text

Mitglied der Helmholtz-Gemeinschaft List of Resources Used • #4: Earth icon by Francesco Paleari from The Noun Project • #4: Einstein icon by Roman Rusinov from The Noun Project • #6: FAIR vector logo from official FAIR website • #6: FAIR rendering from official website • #11: Flare Gun icon by Jop van der Kroef from The Noun Project • #27: STT event animation by Marius C. Mertens • #35: Graphics cards images by NVIDIA promotion • #35: GPU Specifications – Tesla K20X Specifications: http://www.nvidia.com/content/PDF/kepler/Tesla- K20X-BD-06397-001-v07.pdf – Tesla K40 Specifications: http://www.nvidia.com/content/PDF/kepler/Tesla-K40- Active-Board-Spec-BD-06949-001_v03.pdf – Tesla Familiy Overview: http://www.nvidia.com/content/tesla/pdf/NVIDIA-Tesla- Kepler-Family-Datasheet.pdf 38

Slide 72

Slide 72 text

Mitglied der Helmholtz-Gemeinschaft BACKUP 39

Slide 73

Slide 73 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Principle 40 Back

Slide 74

Slide 74 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Principle 40 x y Back

Slide 75

Slide 75 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Principle 40 x y Back

Slide 76

Slide 76 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Principle 40 x y * * (r, α) 1 rij = cos ↵j · xi + sin ↵j · yi + ⇢i Back

Slide 77

Slide 77 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Principle 40 x y * * r α (r, α) 1 rij = cos ↵j · xi + sin ↵j · yi + ⇢i Back

Slide 78

Slide 78 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Principle 40 x y * * r α (r, α) 1 rij = cos ↵j · xi + sin ↵j · yi + ⇢i Back

Slide 79

Slide 79 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Principle 40 x y * * r α (r, α) 1 (r, α)2 rij = cos ↵j · xi + sin ↵j · yi + ⇢i Back

Slide 80

Slide 80 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Principle 40 x y * * r α rij = cos ↵j · xi + sin ↵j · yi + ⇢i Back

Slide 81

Slide 81 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Principle 40 x y * * r α rij = cos ↵j · xi + sin ↵j · yi + ⇢i Back

Slide 82

Slide 82 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Principle 40 x y * * r α rij = cos ↵j · xi + sin ↵j · yi + ⇢i Back

Slide 83

Slide 83 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Principle 40 x y * * r α rij = cos ↵j · xi + sin ↵j · yi + ⇢i Back

Slide 84

Slide 84 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Principle 40 x y * * r α rij = cos ↵j · xi + sin ↵j · yi + ⇢i Back

Slide 85

Slide 85 text

Mitglied der Helmholtz-Gemeinschaft Hough Transform — Principle 40 x y → Bin with highest multiplicity gives track parameters * * r α rij = cos ↵j · xi + sin ↵j · yi + ⇢i Back

Slide 86

Slide 86 text

Mitglied der Helmholtz-Gemeinschaft 41 Riemann Algorithm — Procedure • Create triplet of hit points – All possible three hit combinations need to become triplets 1

Slide 87

Slide 87 text

Mitglied der Helmholtz-Gemeinschaft 41 Riemann Algorithm — Procedure • Create triplet of hit points – All possible three hit combinations need to become triplets • Grow triplets to tracks: Continuously test next hit if it fits to triplet track – Use Riemann paraboloid to circle fit track • Test closeness of new hit: good → add hit; bad → dismiss hit • Continue with next hit – Helix fit: arc length s vs. z position 1 2

Slide 88

Slide 88 text

Mitglied der Helmholtz-Gemeinschaft 42 1 2 3 4 5 1 2 3 4 5 Riemann Track Finder — 1 Seeds 1 Layer number Back

Slide 89

Slide 89 text

Mitglied der Helmholtz-Gemeinschaft 42 1 2 3 4 5 1 2 3 4 5 Riemann Track Finder — 1 Seeds 1 Layer number Back

Slide 90

Slide 90 text

Mitglied der Helmholtz-Gemeinschaft 42 1 2 3 4 5 1 2 3 4 5 Riemann Track Finder — 1 Seeds 1 Layer number Back

Slide 91

Slide 91 text

Mitglied der Helmholtz-Gemeinschaft 42 1 2 3 4 5 21 11 31 1 2 3 4 5 Riemann Track Finder — 1 Seeds 1 Layer number Back

Slide 92

Slide 92 text

Mitglied der Helmholtz-Gemeinschaft 42 1 2 3 4 5 21 11 31 31 11 41 1 2 3 4 5 Riemann Track Finder — 1 Seeds 1 Layer number Back

Slide 93

Slide 93 text

Mitglied der Helmholtz-Gemeinschaft 42 1 2 3 4 5 21 11 31 31 11 41 31 11 32 1 2 3 4 5 Riemann Track Finder — 1 Seeds 1 Layer number Back

Slide 94

Slide 94 text

Mitglied der Helmholtz-Gemeinschaft 42 1 2 3 4 5 21 11 31 31 11 41 31 11 32 1 2 3 4 5 Riemann Track Finder — 1 Seeds 1 Layer number Back

Slide 95

Slide 95 text

Mitglied der Helmholtz-Gemeinschaft 43 Riemann Algorithm — 1 Expansion 2 Back

Slide 96

Slide 96 text

Mitglied der Helmholtz-Gemeinschaft 43 Riemann Algorithm — 1 Expansion 2 x x x x y z‘ Expand to z‘ Back

Slide 97

Slide 97 text

Mitglied der Helmholtz-Gemeinschaft 43 Riemann Algorithm — 1 Expansion 2 x x x x y z‘ Expand to z‘ x x x y x Riemann Surface (paraboloid) Back

Slide 98

Slide 98 text

Mitglied der Helmholtz-Gemeinschaft 43 Riemann Algorithm — 1 Expansion 2 x x x x y z‘ Expand to z‘ x x x y x Riemann Surface (paraboloid) Back

Slide 99

Slide 99 text

Mitglied der Helmholtz-Gemeinschaft 43 Riemann Algorithm — 1 Expansion 2 x x x x y z‘ Expand to z‘ x x x y x Riemann Surface (paraboloid) Back

Slide 100

Slide 100 text

Mitglied der Helmholtz-Gemeinschaft 43 Riemann Algorithm — 1 Expansion 2 x x x x y z‘ Expand to z‘ x x x y x Riemann Surface (paraboloid) Back

Slide 101

Slide 101 text

Mitglied der Helmholtz-Gemeinschaft 43 Riemann Algorithm — 1 Expansion 2 x x x x y z‘ Expand to z‘ x x x y x Riemann Surface (paraboloid) Back

Slide 102

Slide 102 text

Mitglied der Helmholtz-Gemeinschaft 43 Riemann Algorithm — 1 Expansion 2 x x x x y z‘ Expand to z‘ x x x y x Riemann Surface (paraboloid) x Back

Slide 103

Slide 103 text

Mitglied der Helmholtz-Gemeinschaft 43 Riemann Algorithm — 1 Expansion 2 x x x x y z‘ Expand to z‘ x x x y x Riemann Surface (paraboloid) x Back

Slide 104

Slide 104 text

Mitglied der Helmholtz-Gemeinschaft 43 Riemann Algorithm — 1 Expansion 2 x x x x y z‘ Expand to z‘ x x x y x Riemann Surface (paraboloid) x Back

Slide 105

Slide 105 text

Mitglied der Helmholtz-Gemeinschaft 43 Riemann Algorithm — 1 Expansion 2 x x x x y z‘ Expand to z‘ x x x y x Riemann Surface (paraboloid) x Back

Slide 106

Slide 106 text

Mitglied der Helmholtz-Gemeinschaft 43 Riemann Algorithm — 1 Expansion 2 x x x x y z‘ Expand to z‘ x x x y x Riemann Surface (paraboloid) x Back

Slide 107

Slide 107 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Method 44 STT More

Slide 108

Slide 108 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Method 44 STT More

Slide 109

Slide 109 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Method 44 STT More

Slide 110

Slide 110 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Method 44 STT More

Slide 111

Slide 111 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Method • STT hit in pivot straw 44 STT More

Slide 112

Slide 112 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Method • STT hit in pivot straw • Find surrounding hits → Create virtual hit (triplet) at center of gravity (cog) 44 STT More

Slide 113

Slide 113 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Method • STT hit in pivot straw • Find surrounding hits → Create virtual hit (triplet) at center of gravity (cog) • Combine with 44 STT More

Slide 114

Slide 114 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Method • STT hit in pivot straw • Find surrounding hits → Create virtual hit (triplet) at center of gravity (cog) • Combine with 1.Second STT pivot-cog virtual hit 44 STT More

Slide 115

Slide 115 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Method • STT hit in pivot straw • Find surrounding hits → Create virtual hit (triplet) at center of gravity (cog) • Combine with 1.Second STT pivot-cog virtual hit 44 STT More

Slide 116

Slide 116 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Method • STT hit in pivot straw • Find surrounding hits → Create virtual hit (triplet) at center of gravity (cog) • Combine with 1.Second STT pivot-cog virtual hit 44 STT More

Slide 117

Slide 117 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Method • STT hit in pivot straw • Find surrounding hits → Create virtual hit (triplet) at center of gravity (cog) • Combine with 1.Second STT pivot-cog virtual hit 2.Interaction point 44 Interaction Point STT More

Slide 118

Slide 118 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Method • STT hit in pivot straw • Find surrounding hits → Create virtual hit (triplet) at center of gravity (cog) • Combine with 1.Second STT pivot-cog virtual hit 2.Interaction point • Calculate circle through three points 44 Interaction Point STT More

Slide 119

Slide 119 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Method • STT hit in pivot straw • Find surrounding hits → Create virtual hit (triplet) at center of gravity (cog) • Combine with 1.Second STT pivot-cog virtual hit 2.Interaction point • Calculate circle through three points → Track Candidate 44 Interaction Point STT More

Slide 120

Slide 120 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Optimizations 45 • Sector Row testing – After found track: Hit association not with all hits of current window, but only with subset (first test rows of sector, then hits of row) More

Slide 121

Slide 121 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Optimizations 45 • Sector Row testing – After found track: Hit association not with all hits of current window, but only with subset (first test rows of sector, then hits of row) More

Slide 122

Slide 122 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Optimizations 45 • Sector Row testing – After found track: Hit association not with all hits of current window, but only with subset (first test rows of sector, then hits of row) More

Slide 123

Slide 123 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Optimizations 45 • Sector Row testing – After found track: Hit association not with all hits of current window, but only with subset (first test rows of sector, then hits of row) More

Slide 124

Slide 124 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Optimizations 45 • Sector Row testing – After found track: Hit association not with all hits of current window, but only with subset (first test rows of sector, then hits of row) More

Slide 125

Slide 125 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Optimizations • Sector Row testing – Thicken track; shrink sector row layer to line – Find intersection 46 Sector-Row Testing Track Sector-Row Track Sector-Row Back

Slide 126

Slide 126 text

Mitglied der Helmholtz-Gemeinschaft 47 Triplet Finder — Sector Rows

Slide 127

Slide 127 text

Mitglied der Helmholtz-Gemeinschaft Triplet Finder — Kernel Launch Strategies • Joined Kernel (JK): slowest – High # registers → low occupancy • Dynamic Parallelism (DP) / Host Streams (HS): comparable performance – Performance • HS faster for small # processed hits, DP faster for > 45000 hits • HS stagnates there, while DP continues rising – Limiting factor • High # of required kernel calls • Kernel launch latency • Memcopy – HS more affected by this, because • More PCI-E transfers (launch configurations for kernels) • Less launch throughput, kernel launch latency gets more important • False dependencies of launched kernels – Single CPU thread handles all CUDA streams (Multi-thread possible, but synchronization overhead too high for good performance) – Grid scheduling done on hardware (Grid Management Unit) (DP: software) » False dependencies when N(streams) > N(device connections)=323.5 48 Back

Slide 128

Slide 128 text

Mitglied der Helmholtz-Gemeinschaft 49 Triplet Finder — Host Stream Connections Preliminary (in publication)

Slide 129

Slide 129 text

Mitglied der Helmholtz-Gemeinschaft 50 Triplet Finder — Bunch Sizes Preliminary (in publication)