ATLAS Offline Software
Loading...
Searching...
No Matches
GPUClusterInfoAndMomentsCalculatorImplHelper.h
Go to the documentation of this file.
1//
2// Copyright (C) 2002-2025 CERN for the benefit of the ATLAS collaboration
3//
4// Dear emacs, this is -*- c++ -*-
5//
6
7#ifndef CALORECGPU_GPUCLUSTERINFOANDMOMENTSCALCULATORHELPER_CUDA_H
8#define CALORECGPU_GPUCLUSTERINFOANDMOMENTSCALCULATORHELPER_CUDA_H
9
10#include "CaloRecGPU/Helpers.h"
11#include "CaloRecGPU/CUDAFriendlyClasses.h"
12#include "GPUClusterInfoAndMomentsCalculatorImpl.h"
13#include "FPHelpers.h"
14#include "TemporaryHelpers.h"
15
16#include "CaloGeoHelpers/CaloSampling.h"
17//Just enums and stuff, CUDA compatible.
18
19#include <cmath>
20#include <type_traits>
21
22namespace
23{
24 namespace CMCTemporaries
25 {
26 // BIG TABLE OF VARIABLE COEXISTENCE!
27 // I know this is a big comment to have here, but it is the most immediate way to have a reference for what lives where.
28 //
29 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
30 // | CLUSTER STORAGE VARIABLE || Isolation Cluster Pass | Isolation Cell Pass | Zeroth Cluster Pass | First Cell Pass | First Cluster Pass | Second Cell Pass | Shower Axis Pass | Second Cluster Pass | Third Cell Pass | Third Cluster Pass | Final Cluster Pass || Comments |
31 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
32 // | clusterEnergy || | | clusterEnergy | clusterEnergy | clusterEnergy | clusterEnergy | clusterEnergy | clusterEnergy | clusterEnergy | clusterEnergy | clusterEnergy || |
33 // | clusterEt || | | | | clusterEt | clusterEt | clusterEt | clusterEt | clusterEt | clusterEt | clusterEt || Not stored as such by CPU. |
34 // | clusterEta || | | clusterEta | clusterEta | clusterEta | clusterEta | clusterEta | clusterEta | clusterEta | clusterEta | clusterEta || |
35 // | clusterPhi || | | clusterPhi | clusterPhi | clusterPhi | clusterPhi | clusterPhi | clusterPhi | clusterPhi | clusterPhi | clusterPhi || |
36 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
37 // | energyPerSample [00] || energyPerSample [00] | energyPerSample [00] |ePS>absolEnrgPrSmpleAux[00] |AbsolutEnergyPrSmpleAux[00] |AbsolutEnergyPrSmpleAux[00] | | | energyPerSample [00] | energyPerSample [00] | energyPerSample [00] | energyPerSample [00] || |
38 // | energyPerSample [01] || energyPerSample [01] | energyPerSample [01] |ePS>absolEnrgPrSmpleAux[01] |AbsolutEnergyPrSmpleAux[01] |AbsolutEnergyPrSmpleAux[01] | | | energyPerSample [01] | energyPerSample [01] | energyPerSample [01] | energyPerSample [01] || |
39 // | energyPerSample [02] || energyPerSample [02] | energyPerSample [02] |ePS>absolEnrgPrSmpleAux[02] |AbsolutEnergyPrSmpleAux[02] |AbsolutEnergyPrSmpleAux[02] | | | energyPerSample [02] | energyPerSample [02] | energyPerSample [02] | energyPerSample [02] || |
40 // | energyPerSample [03] || energyPerSample [03] | energyPerSample [03] |ePS>absolEnrgPrSmpleAux[03] |AbsolutEnergyPrSmpleAux[03] |AbsolutEnergyPrSmpleAux[03] | | | energyPerSample [03] | energyPerSample [03] | energyPerSample [03] | energyPerSample [03] || |
41 // | energyPerSample [04] || energyPerSample [04] | energyPerSample [04] |ePS>absolEnrgPrSmpleAux[04] |AbsolutEnergyPrSmpleAux[04] |AbsolutEnergyPrSmpleAux[04] | | | energyPerSample [04] | energyPerSample [04] | energyPerSample [04] | energyPerSample [04] || |
42 // | energyPerSample [05] || energyPerSample [05] | energyPerSample [05] |ePS>absolEnrgPrSmpleAux[05] |AbsolutEnergyPrSmpleAux[05] |AbsolutEnergyPrSmpleAux[05] | | | energyPerSample [05] | energyPerSample [05] | energyPerSample [05] | energyPerSample [05] || |
43 // | energyPerSample [06] || energyPerSample [06] | energyPerSample [06] |ePS>absolEnrgPrSmpleAux[06] |AbsolutEnergyPrSmpleAux[06] |AbsolutEnergyPrSmpleAux[06] | | | energyPerSample [06] | energyPerSample [06] | energyPerSample [06] | energyPerSample [06] || |
44 // | energyPerSample [07] || energyPerSample [07] | energyPerSample [07] |ePS>absolEnrgPrSmpleAux[07] |AbsolutEnergyPrSmpleAux[07] |AbsolutEnergyPrSmpleAux[07] | | | energyPerSample [07] | energyPerSample [07] | energyPerSample [07] | energyPerSample [07] || |
45 // | energyPerSample [08] || energyPerSample [08] | energyPerSample [08] |ePS>absolEnrgPrSmpleAux[08] |AbsolutEnergyPrSmpleAux[08] |AbsolutEnergyPrSmpleAux[08] | | | energyPerSample [08] | energyPerSample [08] | energyPerSample [08] | energyPerSample [08] || |
46 // | energyPerSample [09] || energyPerSample [09] | energyPerSample [09] |ePS>absolEnrgPrSmpleAux[09] |AbsolutEnergyPrSmpleAux[09] |AbsolutEnergyPrSmpleAux[09] | | | energyPerSample [09] | energyPerSample [09] | energyPerSample [09] | energyPerSample [09] || |
47 // | energyPerSample [10] || energyPerSample [10] | energyPerSample [10] |ePS>absolEnrgPrSmpleAux[10] |AbsolutEnergyPrSmpleAux[10] |AbsolutEnergyPrSmpleAux[10] | | | energyPerSample [10] | energyPerSample [10] | energyPerSample [10] | energyPerSample [10] || |
48 // | energyPerSample [11] || energyPerSample [11] | energyPerSample [11] |ePS>absolEnrgPrSmpleAux[11] |AbsolutEnergyPrSmpleAux[11] |AbsolutEnergyPrSmpleAux[11] | | | energyPerSample [11] | energyPerSample [11] | energyPerSample [11] | energyPerSample [11] || |
49 // | energyPerSample [12] || energyPerSample [12] | energyPerSample [12] |ePS>absolEnrgPrSmpleAux[12] |AbsolutEnergyPrSmpleAux[12] |AbsolutEnergyPrSmpleAux[12] | | | energyPerSample [12] | energyPerSample [12] | energyPerSample [12] | energyPerSample [12] || |
50 // | energyPerSample [13] || energyPerSample [13] | energyPerSample [13] |ePS>absolEnrgPrSmpleAux[13] |AbsolutEnergyPrSmpleAux[13] |AbsolutEnergyPrSmpleAux[13] | | | energyPerSample [13] | energyPerSample [13] | energyPerSample [13] | energyPerSample [13] || |
51 // | energyPerSample [14] || energyPerSample [14] | energyPerSample [14] |ePS>absolEnrgPrSmpleAux[14] |AbsolutEnergyPrSmpleAux[14] |AbsolutEnergyPrSmpleAux[14] | | | energyPerSample [14] | energyPerSample [14] | energyPerSample [14] | energyPerSample [14] || |
52 // | energyPerSample [15] || energyPerSample [15] | energyPerSample [15] |ePS>absolEnrgPrSmpleAux[15] |AbsolutEnergyPrSmpleAux[15] |AbsolutEnergyPrSmpleAux[15] | | | energyPerSample [15] | energyPerSample [15] | energyPerSample [15] | energyPerSample [15] || |
53 // | energyPerSample [16] || energyPerSample [16] | energyPerSample [16] |ePS>absolEnrgPrSmpleAux[16] |AbsolutEnergyPrSmpleAux[16] |AbsolutEnergyPrSmpleAux[16] | | | energyPerSample [16] | energyPerSample [16] | energyPerSample [16] | energyPerSample [16] || |
54 // | energyPerSample [17] || energyPerSample [17] | energyPerSample [17] |ePS>absolEnrgPrSmpleAux[17] |AbsolutEnergyPrSmpleAux[17] |AbsolutEnergyPrSmpleAux[17] | | | energyPerSample [17] | energyPerSample [17] | energyPerSample [17] | energyPerSample [17] || |
55 // | energyPerSample [18] || energyPerSample [18] | energyPerSample [18] |ePS>absolEnrgPrSmpleAux[18] |AbsolutEnergyPrSmpleAux[18] |AbsolutEnergyPrSmpleAux[18] | | | energyPerSample [18] | energyPerSample [18] | energyPerSample [18] | energyPerSample [18] || |
56 // | energyPerSample [19] || energyPerSample [19] | energyPerSample [19] |ePS>absolEnrgPrSmpleAux[19] |AbsolutEnergyPrSmpleAux[19] |AbsolutEnergyPrSmpleAux[19] | | | energyPerSample [19] | energyPerSample [19] | energyPerSample [19] | energyPerSample [19] || |
57 // | energyPerSample [20] || energyPerSample [20] | energyPerSample [20] |ePS>absolEnrgPrSmpleAux[20] |AbsolutEnergyPrSmpleAux[20] |AbsolutEnergyPrSmpleAux[20] | | | energyPerSample [20] | energyPerSample [20] | energyPerSample [20] | energyPerSample [20] || |
58 // | energyPerSample [21] || energyPerSample [21] | energyPerSample [21] |ePS>absolEnrgPrSmpleAux[21] |AbsolutEnergyPrSmpleAux[21] |AbsolutEnergyPrSmpleAux[21] | | | energyPerSample [21] | energyPerSample [21] | energyPerSample [21] | energyPerSample [21] || |
59 // | energyPerSample [22] || energyPerSample [22] | energyPerSample [22] |ePS>absolEnrgPrSmpleAux[22] |AbsolutEnergyPrSmpleAux[22] |AbsolutEnergyPrSmpleAux[22] | | | energyPerSample [22] | energyPerSample [22] | energyPerSample [22] | energyPerSample [22] || |
60 // | energyPerSample [23] || energyPerSample [23] | energyPerSample [23] |ePS>absolEnrgPrSmpleAux[23] |AbsolutEnergyPrSmpleAux[23] |AbsolutEnergyPrSmpleAux[23] | | | energyPerSample [23] | energyPerSample [23] | energyPerSample [23] | energyPerSample [23] || |
61 // | energyPerSample [24] || energyPerSample [24] | energyPerSample [24] |ePS>absolEnrgPrSmpleAux[24] |AbsolutEnergyPrSmpleAux[24] |AbsolutEnergyPrSmpleAux[24] | | | energyPerSample [24] | energyPerSample [24] | energyPerSample [24] | energyPerSample [24] || |
62 // | energyPerSample [25] || energyPerSample [25] | energyPerSample [25] |ePS>absolEnrgPrSmpleAux[25] |AbsolutEnergyPrSmpleAux[25] |AbsolutEnergyPrSmpleAux[25] | | | energyPerSample [25] | energyPerSample [25] | energyPerSample [25] | energyPerSample [25] || |
63 // | energyPerSample [26] || energyPerSample [26] | energyPerSample [26] |ePS>absolEnrgPrSmpleAux[26] |AbsolutEnergyPrSmpleAux[26] |AbsolutEnergyPrSmpleAux[26] | | | energyPerSample [26] | energyPerSample [26] | energyPerSample [26] | energyPerSample [26] || |
64 // | energyPerSample [27] || energyPerSample [27] | energyPerSample [27] |ePS>absolEnrgPrSmpleAux[27] |AbsolutEnergyPrSmpleAux[27] |AbsolutEnergyPrSmpleAux[27] | | | energyPerSample [27] | energyPerSample [27] | energyPerSample [27] | energyPerSample [27] || |
65 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
66 // | maxEPerSample [00] || energyPerSampleAux [00] | energyPerSampleAux [00] |ePSA>absolutEnrgyPrSmpl[00] | absolutEnergyPerSample[00] | absolutEnergyPerSample[00] | | | energyPerSampleAux [00] | energyPerSampleAux [00] | energyPerSampleAux [00] | maxEPerSample [00] || |
67 // | maxEPerSample [01] || energyPerSampleAux [01] | energyPerSampleAux [01] |ePSA>absolutEnrgyPrSmpl[01] | absolutEnergyPerSample[01] | absolutEnergyPerSample[01] | | | energyPerSampleAux [01] | energyPerSampleAux [01] | energyPerSampleAux [01] | maxEPerSample [01] || |
68 // | maxEPerSample [02] || energyPerSampleAux [02] | energyPerSampleAux [02] |ePSA>absolutEnrgyPrSmpl[02] | absolutEnergyPerSample[02] | absolutEnergyPerSample[02] | | | energyPerSampleAux [02] | energyPerSampleAux [02] | energyPerSampleAux [02] | maxEPerSample [02] || |
69 // | maxEPerSample [03] || energyPerSampleAux [03] | energyPerSampleAux [03] |ePSA>absolutEnrgyPrSmpl[03] | absolutEnergyPerSample[03] | absolutEnergyPerSample[03] | | | energyPerSampleAux [03] | energyPerSampleAux [03] | energyPerSampleAux [03] | maxEPerSample [03] || |
70 // | maxEPerSample [04] || energyPerSampleAux [04] | energyPerSampleAux [04] |ePSA>absolutEnrgyPrSmpl[04] | absolutEnergyPerSample[04] | absolutEnergyPerSample[04] | | | energyPerSampleAux [04] | energyPerSampleAux [04] | energyPerSampleAux [04] | maxEPerSample [04] || |
71 // | maxEPerSample [05] || energyPerSampleAux [05] | energyPerSampleAux [05] |ePSA>absolutEnrgyPrSmpl[05] | absolutEnergyPerSample[05] | absolutEnergyPerSample[05] | | | energyPerSampleAux [05] | energyPerSampleAux [05] | energyPerSampleAux [05] | maxEPerSample [05] || |
72 // | maxEPerSample [06] || energyPerSampleAux [06] | energyPerSampleAux [06] |ePSA>absolutEnrgyPrSmpl[06] | absolutEnergyPerSample[06] | absolutEnergyPerSample[06] | | | energyPerSampleAux [06] | energyPerSampleAux [06] | energyPerSampleAux [06] | maxEPerSample [06] || |
73 // | maxEPerSample [07] || energyPerSampleAux [07] | energyPerSampleAux [07] |ePSA>absolutEnrgyPrSmpl[07] | absolutEnergyPerSample[07] | absolutEnergyPerSample[07] | | | energyPerSampleAux [07] | energyPerSampleAux [07] | energyPerSampleAux [07] | maxEPerSample [07] || |
74 // | maxEPerSample [08] || energyPerSampleAux [08] | energyPerSampleAux [08] |ePSA>absolutEnrgyPrSmpl[08] | absolutEnergyPerSample[08] | absolutEnergyPerSample[08] | | | energyPerSampleAux [08] | energyPerSampleAux [08] | energyPerSampleAux [08] | maxEPerSample [08] || |
75 // | maxEPerSample [09] || energyPerSampleAux [09] | energyPerSampleAux [09] |ePSA>absolutEnrgyPrSmpl[09] | absolutEnergyPerSample[09] | absolutEnergyPerSample[09] | | | energyPerSampleAux [09] | energyPerSampleAux [09] | energyPerSampleAux [09] | maxEPerSample [09] || |
76 // | maxEPerSample [10] || energyPerSampleAux [10] | energyPerSampleAux [10] |ePSA>absolutEnrgyPrSmpl[10] | absolutEnergyPerSample[10] | absolutEnergyPerSample[10] | | | energyPerSampleAux [10] | energyPerSampleAux [10] | energyPerSampleAux [10] | maxEPerSample [10] || |
77 // | maxEPerSample [11] || energyPerSampleAux [11] | energyPerSampleAux [11] |ePSA>absolutEnrgyPrSmpl[11] | absolutEnergyPerSample[11] | absolutEnergyPerSample[11] | | | energyPerSampleAux [11] | energyPerSampleAux [11] | energyPerSampleAux [11] | maxEPerSample [11] || |
78 // | maxEPerSample [12] || energyPerSampleAux [12] | energyPerSampleAux [12] |ePSA>absolutEnrgyPrSmpl[12] | absolutEnergyPerSample[12] | absolutEnergyPerSample[12] | | | energyPerSampleAux [12] | energyPerSampleAux [12] | energyPerSampleAux [12] | maxEPerSample [12] || |
79 // | maxEPerSample [13] || energyPerSampleAux [13] | energyPerSampleAux [13] |ePSA>absolutEnrgyPrSmpl[13] | absolutEnergyPerSample[13] | absolutEnergyPerSample[13] | | | energyPerSampleAux [13] | energyPerSampleAux [13] | energyPerSampleAux [13] | maxEPerSample [13] || |
80 // | maxEPerSample [14] || energyPerSampleAux [14] | energyPerSampleAux [14] |ePSA>absolutEnrgyPrSmpl[14] | absolutEnergyPerSample[14] | absolutEnergyPerSample[14] | | | energyPerSampleAux [14] | energyPerSampleAux [14] | energyPerSampleAux [14] | maxEPerSample [14] || |
81 // | maxEPerSample [15] || energyPerSampleAux [15] | energyPerSampleAux [15] |ePSA>absolutEnrgyPrSmpl[15] | absolutEnergyPerSample[15] | absolutEnergyPerSample[15] | | | energyPerSampleAux [15] | energyPerSampleAux [15] | energyPerSampleAux [15] | maxEPerSample [15] || |
82 // | maxEPerSample [16] || energyPerSampleAux [16] | energyPerSampleAux [16] |ePSA>absolutEnrgyPrSmpl[16] | absolutEnergyPerSample[16] | absolutEnergyPerSample[16] | | | energyPerSampleAux [16] | energyPerSampleAux [16] | energyPerSampleAux [16] | maxEPerSample [16] || |
83 // | maxEPerSample [17] || energyPerSampleAux [17] | energyPerSampleAux [17] |ePSA>absolutEnrgyPrSmpl[17] | absolutEnergyPerSample[17] | absolutEnergyPerSample[17] | | | energyPerSampleAux [17] | energyPerSampleAux [17] | energyPerSampleAux [17] | maxEPerSample [17] || |
84 // | maxEPerSample [18] || energyPerSampleAux [18] | energyPerSampleAux [18] |ePSA>absolutEnrgyPrSmpl[18] | absolutEnergyPerSample[18] | absolutEnergyPerSample[18] | | | energyPerSampleAux [18] | energyPerSampleAux [18] | energyPerSampleAux [18] | maxEPerSample [18] || |
85 // | maxEPerSample [19] || energyPerSampleAux [19] | energyPerSampleAux [19] |ePSA>absolutEnrgyPrSmpl[19] | absolutEnergyPerSample[19] | absolutEnergyPerSample[19] | | | energyPerSampleAux [19] | energyPerSampleAux [19] | energyPerSampleAux [19] | maxEPerSample [19] || |
86 // | maxEPerSample [20] || energyPerSampleAux [20] | energyPerSampleAux [20] |ePSA>absolutEnrgyPrSmpl[20] | absolutEnergyPerSample[20] | absolutEnergyPerSample[20] | | | energyPerSampleAux [20] | energyPerSampleAux [20] | energyPerSampleAux [20] | maxEPerSample [20] || |
87 // | maxEPerSample [21] || energyPerSampleAux [21] | energyPerSampleAux [21] |ePSA>absolutEnrgyPrSmpl[21] | absolutEnergyPerSample[21] | absolutEnergyPerSample[21] | | | energyPerSampleAux [21] | energyPerSampleAux [21] | energyPerSampleAux [21] | maxEPerSample [21] || |
88 // | maxEPerSample [22] || energyPerSampleAux [22] | energyPerSampleAux [22] |ePSA>absolutEnrgyPrSmpl[22] | absolutEnergyPerSample[22] | absolutEnergyPerSample[22] | | | energyPerSampleAux [22] | energyPerSampleAux [22] | energyPerSampleAux [22] | maxEPerSample [22] || |
89 // | maxEPerSample [23] || energyPerSampleAux [23] | energyPerSampleAux [23] |ePSA>absolutEnrgyPrSmpl[23] | absolutEnergyPerSample[23] | absolutEnergyPerSample[23] | | | energyPerSampleAux [23] | energyPerSampleAux [23] | energyPerSampleAux [23] | maxEPerSample [23] || |
90 // | maxEPerSample [24] || energyPerSampleAux [24] | energyPerSampleAux [24] |ePSA>absolutEnrgyPrSmpl[24] | absolutEnergyPerSample[24] | absolutEnergyPerSample[24] | | | energyPerSampleAux [24] | energyPerSampleAux [24] | energyPerSampleAux [24] | maxEPerSample [24] || |
91 // | maxEPerSample [25] || energyPerSampleAux [25] | energyPerSampleAux [25] |ePSA>absolutEnrgyPrSmpl[25] | absolutEnergyPerSample[25] | absolutEnergyPerSample[25] | | | energyPerSampleAux [25] | energyPerSampleAux [25] | energyPerSampleAux [25] | maxEPerSample [25] || |
92 // | maxEPerSample [26] || energyPerSampleAux [26] | energyPerSampleAux [26] |ePSA>absolutEnrgyPrSmpl[26] | absolutEnergyPerSample[26] | absolutEnergyPerSample[26] | | | energyPerSampleAux [26] | energyPerSampleAux [26] | energyPerSampleAux [26] | maxEPerSample [26] || |
93 // | maxEPerSample [27] || energyPerSampleAux [27] | energyPerSampleAux [27] |ePSA>absolutEnrgyPrSmpl[27] | absolutEnergyPerSample[27] | absolutEnergyPerSample[27] | | | energyPerSampleAux [27] | energyPerSampleAux [27] | energyPerSampleAux [27] | maxEPerSample [27] || |
94 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
95 // | maxPhiPerSample [00] ||numberNonEmptySamplings[00] |numberNonEmptySamplings[00] | nNES>phiPerSampleAux [00] | phiPerSampleAux [00] | phiPerSampleAux [00] | | |maxEnrgyAndCellPrSampl2[00] |maxEnrgyAndCellPrSampl2[00] |maxEnrgyAndCellPrSampl2[00] | maxPhiPerSample [00] || |
96 // | maxPhiPerSample [01] ||numberNonEmptySamplings[01] |numberNonEmptySamplings[01] | nNES>phiPerSampleAux [01] | phiPerSampleAux [01] | phiPerSampleAux [01] | | |maxEnrgyAndCellPrSampl2[01] |maxEnrgyAndCellPrSampl2[01] |maxEnrgyAndCellPrSampl2[01] | maxPhiPerSample [01] || |
97 // | maxPhiPerSample [02] ||numberNonEmptySamplings[02] |numberNonEmptySamplings[02] | nNES>phiPerSampleAux [02] | phiPerSampleAux [02] | phiPerSampleAux [02] | | |maxEnrgyAndCellPrSampl2[02] |maxEnrgyAndCellPrSampl2[02] |maxEnrgyAndCellPrSampl2[02] | maxPhiPerSample [02] || |
98 // | maxPhiPerSample [03] ||numberNonEmptySamplings[03] |numberNonEmptySamplings[03] | nNES>phiPerSampleAux [03] | phiPerSampleAux [03] | phiPerSampleAux [03] | | |maxEnrgyAndCellPrSampl2[03] |maxEnrgyAndCellPrSampl2[03] |maxEnrgyAndCellPrSampl2[03] | maxPhiPerSample [03] || |
99 // | maxPhiPerSample [04] ||numberNonEmptySamplings[04] |numberNonEmptySamplings[04] | nNES>phiPerSampleAux [04] | phiPerSampleAux [04] | phiPerSampleAux [04] | | |maxEnrgyAndCellPrSampl2[04] |maxEnrgyAndCellPrSampl2[04] |maxEnrgyAndCellPrSampl2[04] | maxPhiPerSample [04] || |
100 // | maxPhiPerSample [05] ||numberNonEmptySamplings[05] |numberNonEmptySamplings[05] | nNES>phiPerSampleAux [05] | phiPerSampleAux [05] | phiPerSampleAux [05] | | |maxEnrgyAndCellPrSampl2[05] |maxEnrgyAndCellPrSampl2[05] |maxEnrgyAndCellPrSampl2[05] | maxPhiPerSample [05] || |
101 // | maxPhiPerSample [06] ||numberNonEmptySamplings[06] |numberNonEmptySamplings[06] | nNES>phiPerSampleAux [06] | phiPerSampleAux [06] | phiPerSampleAux [06] | | |maxEnrgyAndCellPrSampl2[06] |maxEnrgyAndCellPrSampl2[06] |maxEnrgyAndCellPrSampl2[06] | maxPhiPerSample [06] || |
102 // | maxPhiPerSample [07] ||numberNonEmptySamplings[07] |numberNonEmptySamplings[07] | nNES>phiPerSampleAux [07] | phiPerSampleAux [07] | phiPerSampleAux [07] | | |maxEnrgyAndCellPrSampl2[07] |maxEnrgyAndCellPrSampl2[07] |maxEnrgyAndCellPrSampl2[07] | maxPhiPerSample [07] || |
103 // | maxPhiPerSample [08] ||numberNonEmptySamplings[08] |numberNonEmptySamplings[08] | nNES>phiPerSampleAux [08] | phiPerSampleAux [08] | phiPerSampleAux [08] | | |maxEnrgyAndCellPrSampl2[08] |maxEnrgyAndCellPrSampl2[08] |maxEnrgyAndCellPrSampl2[08] | maxPhiPerSample [08] || |
104 // | maxPhiPerSample [09] ||numberNonEmptySamplings[09] |numberNonEmptySamplings[09] | nNES>phiPerSampleAux [09] | phiPerSampleAux [09] | phiPerSampleAux [09] | | |maxEnrgyAndCellPrSampl2[09] |maxEnrgyAndCellPrSampl2[09] |maxEnrgyAndCellPrSampl2[09] | maxPhiPerSample [09] || |
105 // | maxPhiPerSample [10] ||numberNonEmptySamplings[10] |numberNonEmptySamplings[10] | nNES>phiPerSampleAux [10] | phiPerSampleAux [10] | phiPerSampleAux [10] | | |maxEnrgyAndCellPrSampl2[10] |maxEnrgyAndCellPrSampl2[10] |maxEnrgyAndCellPrSampl2[10] | maxPhiPerSample [10] || |
106 // | maxPhiPerSample [11] ||numberNonEmptySamplings[11] |numberNonEmptySamplings[11] | nNES>phiPerSampleAux [11] | phiPerSampleAux [11] | phiPerSampleAux [11] | | |maxEnrgyAndCellPrSampl2[11] |maxEnrgyAndCellPrSampl2[11] |maxEnrgyAndCellPrSampl2[11] | maxPhiPerSample [11] || |
107 // | maxPhiPerSample [12] ||numberNonEmptySamplings[12] |numberNonEmptySamplings[12] | nNES>phiPerSampleAux [12] | phiPerSampleAux [12] | phiPerSampleAux [12] | | |maxEnrgyAndCellPrSampl2[12] |maxEnrgyAndCellPrSampl2[12] |maxEnrgyAndCellPrSampl2[12] | maxPhiPerSample [12] || |
108 // | maxPhiPerSample [13] ||numberNonEmptySamplings[13] |numberNonEmptySamplings[13] | nNES>phiPerSampleAux [13] | phiPerSampleAux [13] | phiPerSampleAux [13] | | |maxEnrgyAndCellPrSampl2[13] |maxEnrgyAndCellPrSampl2[13] |maxEnrgyAndCellPrSampl2[13] | maxPhiPerSample [13] || |
109 // | maxPhiPerSample [14] ||numberNonEmptySamplings[14] |numberNonEmptySamplings[14] | nNES>phiPerSampleAux [14] | phiPerSampleAux [14] | phiPerSampleAux [14] | | |maxEnrgyAndCellPrSampl2[14] |maxEnrgyAndCellPrSampl2[14] |maxEnrgyAndCellPrSampl2[14] | maxPhiPerSample [14] || |
110 // | maxPhiPerSample [15] ||numberNonEmptySamplings[15] |numberNonEmptySamplings[15] | nNES>phiPerSampleAux [15] | phiPerSampleAux [15] | phiPerSampleAux [15] | | |maxEnrgyAndCellPrSampl2[15] |maxEnrgyAndCellPrSampl2[15] |maxEnrgyAndCellPrSampl2[15] | maxPhiPerSample [15] || |
111 // | maxPhiPerSample [16] ||numberNonEmptySamplings[16] |numberNonEmptySamplings[16] | nNES>phiPerSampleAux [16] | phiPerSampleAux [16] | phiPerSampleAux [16] | | |maxEnrgyAndCellPrSampl2[16] |maxEnrgyAndCellPrSampl2[16] |maxEnrgyAndCellPrSampl2[16] | maxPhiPerSample [16] || |
112 // | maxPhiPerSample [17] ||numberNonEmptySamplings[17] |numberNonEmptySamplings[17] | nNES>phiPerSampleAux [17] | phiPerSampleAux [17] | phiPerSampleAux [17] | | |maxEnrgyAndCellPrSampl2[17] |maxEnrgyAndCellPrSampl2[17] |maxEnrgyAndCellPrSampl2[17] | maxPhiPerSample [17] || |
113 // | maxPhiPerSample [18] ||numberNonEmptySamplings[18] |numberNonEmptySamplings[18] | nNES>phiPerSampleAux [18] | phiPerSampleAux [18] | phiPerSampleAux [18] | | |maxEnrgyAndCellPrSampl2[18] |maxEnrgyAndCellPrSampl2[18] |maxEnrgyAndCellPrSampl2[18] | maxPhiPerSample [18] || |
114 // | maxPhiPerSample [19] ||numberNonEmptySamplings[19] |numberNonEmptySamplings[19] | nNES>phiPerSampleAux [19] | phiPerSampleAux [19] | phiPerSampleAux [19] | | |maxEnrgyAndCellPrSampl2[19] |maxEnrgyAndCellPrSampl2[19] |maxEnrgyAndCellPrSampl2[19] | maxPhiPerSample [19] || |
115 // | maxPhiPerSample [20] ||numberNonEmptySamplings[20] |numberNonEmptySamplings[20] | nNES>phiPerSampleAux [20] | phiPerSampleAux [20] | phiPerSampleAux [20] | | |maxEnrgyAndCellPrSampl2[20] |maxEnrgyAndCellPrSampl2[20] |maxEnrgyAndCellPrSampl2[20] | maxPhiPerSample [20] || |
116 // | maxPhiPerSample [21] ||numberNonEmptySamplings[21] |numberNonEmptySamplings[21] | nNES>phiPerSampleAux [21] | phiPerSampleAux [21] | phiPerSampleAux [21] | | |maxEnrgyAndCellPrSampl2[21] |maxEnrgyAndCellPrSampl2[21] |maxEnrgyAndCellPrSampl2[21] | maxPhiPerSample [21] || |
117 // | maxPhiPerSample [22] ||numberNonEmptySamplings[22] |numberNonEmptySamplings[22] | nNES>phiPerSampleAux [22] | phiPerSampleAux [22] | phiPerSampleAux [22] | | |maxEnrgyAndCellPrSampl2[22] |maxEnrgyAndCellPrSampl2[22] |maxEnrgyAndCellPrSampl2[22] | maxPhiPerSample [22] || |
118 // | maxPhiPerSample [23] ||numberNonEmptySamplings[23] |numberNonEmptySamplings[23] | nNES>phiPerSampleAux [23] | phiPerSampleAux [23] | phiPerSampleAux [23] | | |maxEnrgyAndCellPrSampl2[23] |maxEnrgyAndCellPrSampl2[23] |maxEnrgyAndCellPrSampl2[23] | maxPhiPerSample [23] || |
119 // | maxPhiPerSample [24] ||numberNonEmptySamplings[24] |numberNonEmptySamplings[24] | nNES>phiPerSampleAux [24] | phiPerSampleAux [24] | phiPerSampleAux [24] | | |maxEnrgyAndCellPrSampl2[24] |maxEnrgyAndCellPrSampl2[24] |maxEnrgyAndCellPrSampl2[24] | maxPhiPerSample [24] || |
120 // | maxPhiPerSample [25] ||numberNonEmptySamplings[25] |numberNonEmptySamplings[25] | nNES>phiPerSampleAux [25] | phiPerSampleAux [25] | phiPerSampleAux [25] | | |maxEnrgyAndCellPrSampl2[25] |maxEnrgyAndCellPrSampl2[25] |maxEnrgyAndCellPrSampl2[25] | maxPhiPerSample [25] || |
121 // | maxPhiPerSample [26] ||numberNonEmptySamplings[26] |numberNonEmptySamplings[26] | nNES>phiPerSampleAux [26] | phiPerSampleAux [26] | phiPerSampleAux [26] | | |maxEnrgyAndCellPrSampl2[26] |maxEnrgyAndCellPrSampl2[26] |maxEnrgyAndCellPrSampl2[26] | maxPhiPerSample [26] || |
122 // | maxPhiPerSample [27] ||numberNonEmptySamplings[27] |numberNonEmptySamplings[27] | nNES>phiPerSampleAux [27] | phiPerSampleAux [27] | phiPerSampleAux [27] | | |maxEnrgyAndCellPrSampl2[27] |maxEnrgyAndCellPrSampl2[27] |maxEnrgyAndCellPrSampl2[27] | maxPhiPerSample [27] || |
123 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
124 // | maxEtaPerSample [00] ||maxMomentsEnergPrSample[00] |maxMomentsEnergPrSample[00] | mMEPS>etaPerSampleAux [00] | etaPerSampleAux [00] | etaPerSampleAux [00] | | |maxEnrgyAndCellPrSampl1[00] |maxEnrgyAndCellPrSampl1[00] |maxEnrgyAndCellPrSampl1[00] | maxEtaPerSample [00] || |
125 // | maxEtaPerSample [01] ||maxMomentsEnergPrSample[01] |maxMomentsEnergPrSample[01] | mMEPS>etaPerSampleAux [01] | etaPerSampleAux [01] | etaPerSampleAux [01] | | |maxEnrgyAndCellPrSampl1[01] |maxEnrgyAndCellPrSampl1[01] |maxEnrgyAndCellPrSampl1[01] | maxEtaPerSample [01] || |
126 // | maxEtaPerSample [02] ||maxMomentsEnergPrSample[02] |maxMomentsEnergPrSample[02] | mMEPS>etaPerSampleAux [02] | etaPerSampleAux [02] | etaPerSampleAux [02] | | |maxEnrgyAndCellPrSampl1[02] |maxEnrgyAndCellPrSampl1[02] |maxEnrgyAndCellPrSampl1[02] | maxEtaPerSample [02] || |
127 // | maxEtaPerSample [03] ||maxMomentsEnergPrSample[03] |maxMomentsEnergPrSample[03] | mMEPS>etaPerSampleAux [03] | etaPerSampleAux [03] | etaPerSampleAux [03] | | |maxEnrgyAndCellPrSampl1[03] |maxEnrgyAndCellPrSampl1[03] |maxEnrgyAndCellPrSampl1[03] | maxEtaPerSample [03] || |
128 // | maxEtaPerSample [04] ||maxMomentsEnergPrSample[04] |maxMomentsEnergPrSample[04] | mMEPS>etaPerSampleAux [04] | etaPerSampleAux [04] | etaPerSampleAux [04] | | |maxEnrgyAndCellPrSampl1[04] |maxEnrgyAndCellPrSampl1[04] |maxEnrgyAndCellPrSampl1[04] | maxEtaPerSample [04] || |
129 // | maxEtaPerSample [05] ||maxMomentsEnergPrSample[05] |maxMomentsEnergPrSample[05] | mMEPS>etaPerSampleAux [05] | etaPerSampleAux [05] | etaPerSampleAux [05] | | |maxEnrgyAndCellPrSampl1[05] |maxEnrgyAndCellPrSampl1[05] |maxEnrgyAndCellPrSampl1[05] | maxEtaPerSample [05] || |
130 // | maxEtaPerSample [06] ||maxMomentsEnergPrSample[06] |maxMomentsEnergPrSample[06] | mMEPS>etaPerSampleAux [06] | etaPerSampleAux [06] | etaPerSampleAux [06] | | |maxEnrgyAndCellPrSampl1[06] |maxEnrgyAndCellPrSampl1[06] |maxEnrgyAndCellPrSampl1[06] | maxEtaPerSample [06] || |
131 // | maxEtaPerSample [07] ||maxMomentsEnergPrSample[07] |maxMomentsEnergPrSample[07] | mMEPS>etaPerSampleAux [07] | etaPerSampleAux [07] | etaPerSampleAux [07] | | |maxEnrgyAndCellPrSampl1[07] |maxEnrgyAndCellPrSampl1[07] |maxEnrgyAndCellPrSampl1[07] | maxEtaPerSample [07] || |
132 // | maxEtaPerSample [08] ||maxMomentsEnergPrSample[08] |maxMomentsEnergPrSample[08] | mMEPS>etaPerSampleAux [08] | etaPerSampleAux [08] | etaPerSampleAux [08] | | |maxEnrgyAndCellPrSampl1[08] |maxEnrgyAndCellPrSampl1[08] |maxEnrgyAndCellPrSampl1[08] | maxEtaPerSample [08] || |
133 // | maxEtaPerSample [09] ||maxMomentsEnergPrSample[09] |maxMomentsEnergPrSample[09] | mMEPS>etaPerSampleAux [09] | etaPerSampleAux [09] | etaPerSampleAux [09] | | |maxEnrgyAndCellPrSampl1[09] |maxEnrgyAndCellPrSampl1[09] |maxEnrgyAndCellPrSampl1[09] | maxEtaPerSample [09] || |
134 // | maxEtaPerSample [10] ||maxMomentsEnergPrSample[10] |maxMomentsEnergPrSample[10] | mMEPS>etaPerSampleAux [10] | etaPerSampleAux [10] | etaPerSampleAux [10] | | |maxEnrgyAndCellPrSampl1[10] |maxEnrgyAndCellPrSampl1[10] |maxEnrgyAndCellPrSampl1[10] | maxEtaPerSample [10] || |
135 // | maxEtaPerSample [11] ||maxMomentsEnergPrSample[11] |maxMomentsEnergPrSample[11] | mMEPS>etaPerSampleAux [11] | etaPerSampleAux [11] | etaPerSampleAux [11] | | |maxEnrgyAndCellPrSampl1[11] |maxEnrgyAndCellPrSampl1[11] |maxEnrgyAndCellPrSampl1[11] | maxEtaPerSample [11] || |
136 // | maxEtaPerSample [12] ||maxMomentsEnergPrSample[12] |maxMomentsEnergPrSample[12] | mMEPS>etaPerSampleAux [12] | etaPerSampleAux [12] | etaPerSampleAux [12] | | |maxEnrgyAndCellPrSampl1[12] |maxEnrgyAndCellPrSampl1[12] |maxEnrgyAndCellPrSampl1[12] | maxEtaPerSample [12] || |
137 // | maxEtaPerSample [13] ||maxMomentsEnergPrSample[13] |maxMomentsEnergPrSample[13] | mMEPS>etaPerSampleAux [13] | etaPerSampleAux [13] | etaPerSampleAux [13] | | |maxEnrgyAndCellPrSampl1[13] |maxEnrgyAndCellPrSampl1[13] |maxEnrgyAndCellPrSampl1[13] | maxEtaPerSample [13] || |
138 // | maxEtaPerSample [14] ||maxMomentsEnergPrSample[14] |maxMomentsEnergPrSample[14] | mMEPS>etaPerSampleAux [14] | etaPerSampleAux [14] | etaPerSampleAux [14] | | |maxEnrgyAndCellPrSampl1[14] |maxEnrgyAndCellPrSampl1[14] |maxEnrgyAndCellPrSampl1[14] | maxEtaPerSample [14] || |
139 // | maxEtaPerSample [15] ||maxMomentsEnergPrSample[15] |maxMomentsEnergPrSample[15] | mMEPS>etaPerSampleAux [15] | etaPerSampleAux [15] | etaPerSampleAux [15] | | |maxEnrgyAndCellPrSampl1[15] |maxEnrgyAndCellPrSampl1[15] |maxEnrgyAndCellPrSampl1[15] | maxEtaPerSample [15] || |
140 // | maxEtaPerSample [16] ||maxMomentsEnergPrSample[16] |maxMomentsEnergPrSample[16] | mMEPS>etaPerSampleAux [16] | etaPerSampleAux [16] | etaPerSampleAux [16] | | |maxEnrgyAndCellPrSampl1[16] |maxEnrgyAndCellPrSampl1[16] |maxEnrgyAndCellPrSampl1[16] | maxEtaPerSample [16] || |
141 // | maxEtaPerSample [17] ||maxMomentsEnergPrSample[17] |maxMomentsEnergPrSample[17] | mMEPS>etaPerSampleAux [17] | etaPerSampleAux [17] | etaPerSampleAux [17] | | |maxEnrgyAndCellPrSampl1[17] |maxEnrgyAndCellPrSampl1[17] |maxEnrgyAndCellPrSampl1[17] | maxEtaPerSample [17] || |
142 // | maxEtaPerSample [18] ||maxMomentsEnergPrSample[18] |maxMomentsEnergPrSample[18] | mMEPS>etaPerSampleAux [18] | etaPerSampleAux [18] | etaPerSampleAux [18] | | |maxEnrgyAndCellPrSampl1[18] |maxEnrgyAndCellPrSampl1[18] |maxEnrgyAndCellPrSampl1[18] | maxEtaPerSample [18] || |
143 // | maxEtaPerSample [19] ||maxMomentsEnergPrSample[19] |maxMomentsEnergPrSample[19] | mMEPS>etaPerSampleAux [19] | etaPerSampleAux [19] | etaPerSampleAux [19] | | |maxEnrgyAndCellPrSampl1[19] |maxEnrgyAndCellPrSampl1[19] |maxEnrgyAndCellPrSampl1[19] | maxEtaPerSample [19] || |
144 // | maxEtaPerSample [20] ||maxMomentsEnergPrSample[20] |maxMomentsEnergPrSample[20] | mMEPS>etaPerSampleAux [20] | etaPerSampleAux [20] | etaPerSampleAux [20] | | |maxEnrgyAndCellPrSampl1[20] |maxEnrgyAndCellPrSampl1[20] |maxEnrgyAndCellPrSampl1[20] | maxEtaPerSample [20] || |
145 // | maxEtaPerSample [21] ||maxMomentsEnergPrSample[21] |maxMomentsEnergPrSample[21] | mMEPS>etaPerSampleAux [21] | etaPerSampleAux [21] | etaPerSampleAux [21] | | |maxEnrgyAndCellPrSampl1[21] |maxEnrgyAndCellPrSampl1[21] |maxEnrgyAndCellPrSampl1[21] | maxEtaPerSample [21] || |
146 // | maxEtaPerSample [22] ||maxMomentsEnergPrSample[22] |maxMomentsEnergPrSample[22] | mMEPS>etaPerSampleAux [22] | etaPerSampleAux [22] | etaPerSampleAux [22] | | |maxEnrgyAndCellPrSampl1[22] |maxEnrgyAndCellPrSampl1[22] |maxEnrgyAndCellPrSampl1[22] | maxEtaPerSample [22] || |
147 // | maxEtaPerSample [23] ||maxMomentsEnergPrSample[23] |maxMomentsEnergPrSample[23] | mMEPS>etaPerSampleAux [23] | etaPerSampleAux [23] | etaPerSampleAux [23] | | |maxEnrgyAndCellPrSampl1[23] |maxEnrgyAndCellPrSampl1[23] |maxEnrgyAndCellPrSampl1[23] | maxEtaPerSample [23] || |
148 // | maxEtaPerSample [24] ||maxMomentsEnergPrSample[24] |maxMomentsEnergPrSample[24] | mMEPS>etaPerSampleAux [24] | etaPerSampleAux [24] | etaPerSampleAux [24] | | |maxEnrgyAndCellPrSampl1[24] |maxEnrgyAndCellPrSampl1[24] |maxEnrgyAndCellPrSampl1[24] | maxEtaPerSample [24] || |
149 // | maxEtaPerSample [25] ||maxMomentsEnergPrSample[25] |maxMomentsEnergPrSample[25] | mMEPS>etaPerSampleAux [25] | etaPerSampleAux [25] | etaPerSampleAux [25] | | |maxEnrgyAndCellPrSampl1[25] |maxEnrgyAndCellPrSampl1[25] |maxEnrgyAndCellPrSampl1[25] | maxEtaPerSample [25] || |
150 // | maxEtaPerSample [26] ||maxMomentsEnergPrSample[26] |maxMomentsEnergPrSample[26] | mMEPS>etaPerSampleAux [26] | etaPerSampleAux [26] | etaPerSampleAux [26] | | |maxEnrgyAndCellPrSampl1[26] |maxEnrgyAndCellPrSampl1[26] |maxEnrgyAndCellPrSampl1[26] | maxEtaPerSample [26] || |
151 // | maxEtaPerSample [27] ||maxMomentsEnergPrSample[27] |maxMomentsEnergPrSample[27] | mMEPS>etaPerSampleAux [27] | etaPerSampleAux [27] | etaPerSampleAux [27] | | |maxEnrgyAndCellPrSampl1[27] |maxEnrgyAndCellPrSampl1[27] |maxEnrgyAndCellPrSampl1[27] | maxEtaPerSample [27] || |
152 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
153 // | etaPerSample [00] || numberEmptySamplings [00] | numberEmptySamplings [00] | nuES>etaPerSample [00] | etaPerSample [00] | etaPerSample [00] | etaPerSample [00] | etaPerSample [00] | etaPerSample [00] | etaPerSample [00] | etaPerSample [00] | etaPerSample [00] || |
154 // | etaPerSample [01] || numberEmptySamplings [01] | numberEmptySamplings [01] | nuES>etaPerSample [01] | etaPerSample [01] | etaPerSample [01] | etaPerSample [01] | etaPerSample [01] | etaPerSample [01] | etaPerSample [01] | etaPerSample [01] | etaPerSample [01] || |
155 // | etaPerSample [02] || numberEmptySamplings [02] | numberEmptySamplings [02] | nuES>etaPerSample [02] | etaPerSample [02] | etaPerSample [02] | etaPerSample [02] | etaPerSample [02] | etaPerSample [02] | etaPerSample [02] | etaPerSample [02] | etaPerSample [02] || |
156 // | etaPerSample [03] || numberEmptySamplings [03] | numberEmptySamplings [03] | nuES>etaPerSample [03] | etaPerSample [03] | etaPerSample [03] | etaPerSample [03] | etaPerSample [03] | etaPerSample [03] | etaPerSample [03] | etaPerSample [03] | etaPerSample [03] || |
157 // | etaPerSample [04] || numberEmptySamplings [04] | numberEmptySamplings [04] | nuES>etaPerSample [04] | etaPerSample [04] | etaPerSample [04] | etaPerSample [04] | etaPerSample [04] | etaPerSample [04] | etaPerSample [04] | etaPerSample [04] | etaPerSample [04] || |
158 // | etaPerSample [05] || numberEmptySamplings [05] | numberEmptySamplings [05] | nuES>etaPerSample [05] | etaPerSample [05] | etaPerSample [05] | etaPerSample [05] | etaPerSample [05] | etaPerSample [05] | etaPerSample [05] | etaPerSample [05] | etaPerSample [05] || |
159 // | etaPerSample [06] || numberEmptySamplings [06] | numberEmptySamplings [06] | nuES>etaPerSample [06] | etaPerSample [06] | etaPerSample [06] | etaPerSample [06] | etaPerSample [06] | etaPerSample [06] | etaPerSample [06] | etaPerSample [06] | etaPerSample [06] || |
160 // | etaPerSample [07] || numberEmptySamplings [07] | numberEmptySamplings [07] | nuES>etaPerSample [07] | etaPerSample [07] | etaPerSample [07] | etaPerSample [07] | etaPerSample [07] | etaPerSample [07] | etaPerSample [07] | etaPerSample [07] | etaPerSample [07] || |
161 // | etaPerSample [08] || numberEmptySamplings [08] | numberEmptySamplings [08] | nuES>etaPerSample [08] | etaPerSample [08] | etaPerSample [08] | etaPerSample [08] | etaPerSample [08] | etaPerSample [08] | etaPerSample [08] | etaPerSample [08] | etaPerSample [08] || |
162 // | etaPerSample [09] || numberEmptySamplings [09] | numberEmptySamplings [09] | nuES>etaPerSample [09] | etaPerSample [09] | etaPerSample [09] | etaPerSample [09] | etaPerSample [09] | etaPerSample [09] | etaPerSample [09] | etaPerSample [09] | etaPerSample [09] || |
163 // | etaPerSample [10] || numberEmptySamplings [10] | numberEmptySamplings [10] | nuES>etaPerSample [10] | etaPerSample [10] | etaPerSample [10] | etaPerSample [10] | etaPerSample [10] | etaPerSample [10] | etaPerSample [10] | etaPerSample [10] | etaPerSample [10] || |
164 // | etaPerSample [11] || numberEmptySamplings [11] | numberEmptySamplings [11] | nuES>etaPerSample [11] | etaPerSample [11] | etaPerSample [11] | etaPerSample [11] | etaPerSample [11] | etaPerSample [11] | etaPerSample [11] | etaPerSample [11] | etaPerSample [11] || |
165 // | etaPerSample [12] || numberEmptySamplings [12] | numberEmptySamplings [12] | nuES>etaPerSample [12] | etaPerSample [12] | etaPerSample [12] | etaPerSample [12] | etaPerSample [12] | etaPerSample [12] | etaPerSample [12] | etaPerSample [12] | etaPerSample [12] || |
166 // | etaPerSample [13] || numberEmptySamplings [13] | numberEmptySamplings [13] | nuES>etaPerSample [13] | etaPerSample [13] | etaPerSample [13] | etaPerSample [13] | etaPerSample [13] | etaPerSample [13] | etaPerSample [13] | etaPerSample [13] | etaPerSample [13] || |
167 // | etaPerSample [14] || numberEmptySamplings [14] | numberEmptySamplings [14] | nuES>etaPerSample [14] | etaPerSample [14] | etaPerSample [14] | etaPerSample [14] | etaPerSample [14] | etaPerSample [14] | etaPerSample [14] | etaPerSample [14] | etaPerSample [14] || |
168 // | etaPerSample [15] || numberEmptySamplings [15] | numberEmptySamplings [15] | nuES>etaPerSample [15] | etaPerSample [15] | etaPerSample [15] | etaPerSample [15] | etaPerSample [15] | etaPerSample [15] | etaPerSample [15] | etaPerSample [15] | etaPerSample [15] || |
169 // | etaPerSample [16] || numberEmptySamplings [16] | numberEmptySamplings [16] | nuES>etaPerSample [16] | etaPerSample [16] | etaPerSample [16] | etaPerSample [16] | etaPerSample [16] | etaPerSample [16] | etaPerSample [16] | etaPerSample [16] | etaPerSample [16] || |
170 // | etaPerSample [17] || numberEmptySamplings [17] | numberEmptySamplings [17] | nuES>etaPerSample [17] | etaPerSample [17] | etaPerSample [17] | etaPerSample [17] | etaPerSample [17] | etaPerSample [17] | etaPerSample [17] | etaPerSample [17] | etaPerSample [17] || |
171 // | etaPerSample [18] || numberEmptySamplings [18] | numberEmptySamplings [18] | nuES>etaPerSample [18] | etaPerSample [18] | etaPerSample [18] | etaPerSample [18] | etaPerSample [18] | etaPerSample [18] | etaPerSample [18] | etaPerSample [18] | etaPerSample [18] || |
172 // | etaPerSample [19] || numberEmptySamplings [19] | numberEmptySamplings [19] | nuES>etaPerSample [19] | etaPerSample [19] | etaPerSample [19] | etaPerSample [19] | etaPerSample [19] | etaPerSample [19] | etaPerSample [19] | etaPerSample [19] | etaPerSample [19] || |
173 // | etaPerSample [20] || numberEmptySamplings [20] | numberEmptySamplings [20] | nuES>etaPerSample [20] | etaPerSample [20] | etaPerSample [20] | etaPerSample [20] | etaPerSample [20] | etaPerSample [20] | etaPerSample [20] | etaPerSample [20] | etaPerSample [20] || |
174 // | etaPerSample [21] || numberEmptySamplings [21] | numberEmptySamplings [21] | nuES>etaPerSample [21] | etaPerSample [21] | etaPerSample [21] | etaPerSample [21] | etaPerSample [21] | etaPerSample [21] | etaPerSample [21] | etaPerSample [21] | etaPerSample [21] || |
175 // | etaPerSample [22] || numberEmptySamplings [22] | numberEmptySamplings [22] | nuES>etaPerSample [22] | etaPerSample [22] | etaPerSample [22] | etaPerSample [22] | etaPerSample [22] | etaPerSample [22] | etaPerSample [22] | etaPerSample [22] | etaPerSample [22] || |
176 // | etaPerSample [23] || numberEmptySamplings [23] | numberEmptySamplings [23] | nuES>etaPerSample [23] | etaPerSample [23] | etaPerSample [23] | etaPerSample [23] | etaPerSample [23] | etaPerSample [23] | etaPerSample [23] | etaPerSample [23] | etaPerSample [23] || |
177 // | etaPerSample [24] || numberEmptySamplings [24] | numberEmptySamplings [24] | nuES>etaPerSample [24] | etaPerSample [24] | etaPerSample [24] | etaPerSample [24] | etaPerSample [24] | etaPerSample [24] | etaPerSample [24] | etaPerSample [24] | etaPerSample [24] || |
178 // | etaPerSample [25] || numberEmptySamplings [25] | numberEmptySamplings [25] | nuES>etaPerSample [25] | etaPerSample [25] | etaPerSample [25] | etaPerSample [25] | etaPerSample [25] | etaPerSample [25] | etaPerSample [25] | etaPerSample [25] | etaPerSample [25] || |
179 // | etaPerSample [26] || numberEmptySamplings [26] | numberEmptySamplings [26] | nuES>etaPerSample [26] | etaPerSample [26] | etaPerSample [26] | etaPerSample [26] | etaPerSample [26] | etaPerSample [26] | etaPerSample [26] | etaPerSample [26] | etaPerSample [26] || |
180 // | etaPerSample [27] || numberEmptySamplings [27] | numberEmptySamplings [27] | nuES>etaPerSample [27] | etaPerSample [27] | etaPerSample [27] | etaPerSample [27] | etaPerSample [27] | etaPerSample [27] | etaPerSample [27] | etaPerSample [27] | etaPerSample [27] || |
181 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
182 // | phiPerSample [00] || | | phiPerSample [00] | phiPerSample [00] | phiPerSample [00] | phiPerSample [00] | phiPerSample [00] | phiPerSample [00] | phiPerSample [00] | phiPerSample [00] | phiPerSample [00] || |
183 // | phiPerSample [01] || | | phiPerSample [01] | phiPerSample [01] | phiPerSample [01] | phiPerSample [01] | phiPerSample [01] | phiPerSample [01] | phiPerSample [01] | phiPerSample [01] | phiPerSample [01] || |
184 // | phiPerSample [02] || | | phiPerSample [02] | phiPerSample [02] | phiPerSample [02] | phiPerSample [02] | phiPerSample [02] | phiPerSample [02] | phiPerSample [02] | phiPerSample [02] | phiPerSample [02] || |
185 // | phiPerSample [03] || | | phiPerSample [03] | phiPerSample [03] | phiPerSample [03] | phiPerSample [03] | phiPerSample [03] | phiPerSample [03] | phiPerSample [03] | phiPerSample [03] | phiPerSample [03] || |
186 // | phiPerSample [04] || | | phiPerSample [04] | phiPerSample [04] | phiPerSample [04] | phiPerSample [04] | phiPerSample [04] | phiPerSample [04] | phiPerSample [04] | phiPerSample [04] | phiPerSample [04] || |
187 // | phiPerSample [05] || | | phiPerSample [05] | phiPerSample [05] | phiPerSample [05] | phiPerSample [05] | phiPerSample [05] | phiPerSample [05] | phiPerSample [05] | phiPerSample [05] | phiPerSample [05] || |
188 // | phiPerSample [06] || | | phiPerSample [06] | phiPerSample [06] | phiPerSample [06] | phiPerSample [06] | phiPerSample [06] | phiPerSample [06] | phiPerSample [06] | phiPerSample [06] | phiPerSample [06] || |
189 // | phiPerSample [07] || | | phiPerSample [07] | phiPerSample [07] | phiPerSample [07] | phiPerSample [07] | phiPerSample [07] | phiPerSample [07] | phiPerSample [07] | phiPerSample [07] | phiPerSample [07] || |
190 // | phiPerSample [08] || | | phiPerSample [08] | phiPerSample [08] | phiPerSample [08] | phiPerSample [08] | phiPerSample [08] | phiPerSample [08] | phiPerSample [08] | phiPerSample [08] | phiPerSample [08] || |
191 // | phiPerSample [09] || | | phiPerSample [09] | phiPerSample [09] | phiPerSample [09] | phiPerSample [09] | phiPerSample [09] | phiPerSample [09] | phiPerSample [09] | phiPerSample [09] | phiPerSample [09] || |
192 // | phiPerSample [10] || | | phiPerSample [10] | phiPerSample [10] | phiPerSample [10] | phiPerSample [10] | phiPerSample [10] | phiPerSample [10] | phiPerSample [10] | phiPerSample [10] | phiPerSample [10] || |
193 // | phiPerSample [11] || | | phiPerSample [11] | phiPerSample [11] | phiPerSample [11] | phiPerSample [11] | phiPerSample [11] | phiPerSample [11] | phiPerSample [11] | phiPerSample [11] | phiPerSample [11] || |
194 // | phiPerSample [12] || | | phiPerSample [12] | phiPerSample [12] | phiPerSample [12] | phiPerSample [12] | phiPerSample [12] | phiPerSample [12] | phiPerSample [12] | phiPerSample [12] | phiPerSample [12] || |
195 // | phiPerSample [13] || | | phiPerSample [13] | phiPerSample [13] | phiPerSample [13] | phiPerSample [13] | phiPerSample [13] | phiPerSample [13] | phiPerSample [13] | phiPerSample [13] | phiPerSample [13] || |
196 // | phiPerSample [14] || | | phiPerSample [14] | phiPerSample [14] | phiPerSample [14] | phiPerSample [14] | phiPerSample [14] | phiPerSample [14] | phiPerSample [14] | phiPerSample [14] | phiPerSample [14] || |
197 // | phiPerSample [15] || | | phiPerSample [15] | phiPerSample [15] | phiPerSample [15] | phiPerSample [15] | phiPerSample [15] | phiPerSample [15] | phiPerSample [15] | phiPerSample [15] | phiPerSample [15] || |
198 // | phiPerSample [16] || | | phiPerSample [16] | phiPerSample [16] | phiPerSample [16] | phiPerSample [16] | phiPerSample [16] | phiPerSample [16] | phiPerSample [16] | phiPerSample [16] | phiPerSample [16] || |
199 // | phiPerSample [17] || | | phiPerSample [17] | phiPerSample [17] | phiPerSample [17] | phiPerSample [17] | phiPerSample [17] | phiPerSample [17] | phiPerSample [17] | phiPerSample [17] | phiPerSample [17] || |
200 // | phiPerSample [18] || | | phiPerSample [18] | phiPerSample [18] | phiPerSample [18] | phiPerSample [18] | phiPerSample [18] | phiPerSample [18] | phiPerSample [18] | phiPerSample [18] | phiPerSample [18] || |
201 // | phiPerSample [19] || | | phiPerSample [19] | phiPerSample [19] | phiPerSample [19] | phiPerSample [19] | phiPerSample [19] | phiPerSample [19] | phiPerSample [19] | phiPerSample [19] | phiPerSample [19] || |
202 // | phiPerSample [20] || | | phiPerSample [20] | phiPerSample [20] | phiPerSample [20] | phiPerSample [20] | phiPerSample [20] | phiPerSample [20] | phiPerSample [20] | phiPerSample [20] | phiPerSample [20] || |
203 // | phiPerSample [21] || | | phiPerSample [21] | phiPerSample [21] | phiPerSample [21] | phiPerSample [21] | phiPerSample [21] | phiPerSample [21] | phiPerSample [21] | phiPerSample [21] | phiPerSample [21] || |
204 // | phiPerSample [22] || | | phiPerSample [22] | phiPerSample [22] | phiPerSample [22] | phiPerSample [22] | phiPerSample [22] | phiPerSample [22] | phiPerSample [22] | phiPerSample [22] | phiPerSample [22] || |
205 // | phiPerSample [23] || | | phiPerSample [23] | phiPerSample [23] | phiPerSample [23] | phiPerSample [23] | phiPerSample [23] | phiPerSample [23] | phiPerSample [23] | phiPerSample [23] | phiPerSample [23] || |
206 // | phiPerSample [24] || | | phiPerSample [24] | phiPerSample [24] | phiPerSample [24] | phiPerSample [24] | phiPerSample [24] | phiPerSample [24] | phiPerSample [24] | phiPerSample [24] | phiPerSample [24] || |
207 // | phiPerSample [25] || | | phiPerSample [25] | phiPerSample [25] | phiPerSample [25] | phiPerSample [25] | phiPerSample [25] | phiPerSample [25] | phiPerSample [25] | phiPerSample [25] | phiPerSample [25] || |
208 // | phiPerSample [26] || | | phiPerSample [26] | phiPerSample [26] | phiPerSample [26] | phiPerSample [26] | phiPerSample [26] | phiPerSample [26] | phiPerSample [26] | phiPerSample [26] | phiPerSample [26] || |
209 // | phiPerSample [27] || | | phiPerSample [27] | phiPerSample [27] | phiPerSample [27] | phiPerSample [27] | phiPerSample [27] | phiPerSample [27] | phiPerSample [27] | phiPerSample [27] | phiPerSample [27] || |
210 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
211 // | time || | | | | time | time | time | time | time | time | time || |
212 // | firstPhi || | | firstPhi | firstPhi | firstPhi | firstPhi | firstPhi | firstPhi | firstPhi | firstPhi | firstPhi || |
213 // | firstEta || | | firstEta | firstEta | firstEta | firstEta | firstEta | firstEta | firstEta | firstEta | firstEta || |
214 // | secondR || | | firstPhiAux | firstPhiAux | firstPhiAux -> matrix11 | matrix11 | matrix11 | secondR | secondR | secondR | secondR || |
215 // | secondLambda || | |energyDensityNormlizationAux|energyDensityNormlizationAux|EDnstNormAux>avgLArQNormAux | averageLArQNormAux | averageLArQNormAux |avgLArQNormAux>secondLambda | secondLambda | secondLambda | secondLambda || |
216 // | deltaPhi || | | | | matrix10Aux | matrix10Aux | matrix10Aux -> deltaPhi | deltaPhi | deltaPhi | deltaPhi | deltaPhi || |
217 // | deltaTheta || | | | | matrix20Aux | matrix20Aux | matrix20Aux -> deltaTheta | deltaTheta | deltaTheta | deltaTheta | deltaTheta || |
218 // | deltaAlpha || | | | | matrix21Aux | matrix21Aux | matrix21Aux -> deltaAlpha | deltaAlpha | deltaAlpha | deltaAlpha | deltaAlpha || |
219 // | centerX || | | centerX | centerX | centerX | centerX | centerX | centerX | centerX | centerX | centerX || |
220 // | centerY || | | centerY | centerY | centerY | centerY | centerY | centerY | centerY | centerY | centerY || |
221 // | centerZ || | | centerZ | centerZ | centerZ | centerZ | centerZ | centerZ | centerZ | centerZ | centerZ || |
222 // | centerMag || | | engFracEMAux | engFracEMAux | engFracEMAux | | centerMag | centerMag | centerMag | centerMag | centerMag || |
223 // | centerLambda || | | | | badCellsCorrEAux | badCellsCorrEAux | badCellsCorrEAux | badCellsCorrEAux | | | centerLambda || |
224 // | lateral || | | secondEngDensAux | secondEngDensAux | secondEngDensAux | | | lateral | lateral | lateral | lateral || |
225 // | longitudinal || | | firstEtaAux | firstEtaAux | firstEtaAux -> matrix00 | matrix00 | matrix00 | longitudinal | longitudinal | longitudinal | longitudinal || |
226 // | engFracEM || | | engFracEM | engFracEM | engFracEM | engFracEM | engFracEM | engFracEM | engFracEM | engFracEM | engFracEM || |
227 // | engFracMax || | | | | engFracMax | engFracMax | engFracMax | engFracMax | engFracMax | engFracMax | engFracMax || |
228 // | engFracCore || engFracCore | engFracCore | engFracCore | engFracCore | engFracCore | engFracCore | engFracCore | engFracCore | engFracCore | engFracCore | engFracCore || |
229 // | firstEngDens || | | firstEngDens | firstEngDens | firstEngDens | firstEngDens | firstEngDens | firstEngDens | firstEngDens | firstEngDens | firstEngDens || |
230 // | secondEngDens || | | secondEngDens | secondEngDens | secondEngDens | secondEngDens | secondEngDens | secondEngDens | secondEngDens | secondEngDens | secondEngDens || |
231 // | isolation || | | isolation | isolation | isolation | isolation | isolation | isolation | isolation | isolation | isolation || |
232 // | engBadCells || | | | | engBadCells | engBadCells | engBadCells | engBadCells | engBadCells | engBadCells | engBadCells || |
233 // | nBadCells || | | | | nBadCells | nBadCells | nBadCells | nBadCells | nBadCells | nBadCells | nBadCells || |
234 // | nBadCellsCorr || | | | | nBadCellsCorr | nBadCellsCorr | nBadCellsCorr | nBadCellsCorr | nBadCellsCorr | nBadCellsCorr | nBadCellsCorr || |
235 // | badCellsCorrE || | | | | badCellsCorrE | badCellsCorrE | badCellsCorrE | badCellsCorrE | badCellsCorrE | badCellsCorrE | badCellsCorrE || |
236 // | badLArQFrac || | | | | badLArQFrac | badLArQFrac | badLArQFrac | badLArQFrac | badLArQFrac | badLArQFrac | badLArQFrac || |
237 // | engPos || engPos | engPos | engPos | engPos | engPos | engPos | engPos | engPos | engPos | engPos | engPos || |
238 // | significance || | | | | significance | significance | significance | significance | significance | significance | significance || |
239 // | cellSignificance || | | mYAux | mYAux | mYAux | | | cellSignificance | cellSignificance | cellSignificance | cellSignificance || |
240 // | cellSigSampling || | | mZAux | mZAux | mZAux | | | cellSigSampling | cellSigSampling | cellSigSampling | cellSigSampling || |
241 // | avgLArQ || | | | | avgLArQ | avgLArQ | avgLArQ | avgLArQ | avgLArQ | avgLArQ | avgLArQ || |
242 // | avgTileQ || | | | | avgTileQ | avgTileQ | avgTileQ | avgTileQ | avgTileQ | avgTileQ | avgTileQ || |
243 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
244 // | engBadHVCells || | engPosAux | engPosAux | | timeNormalizationAux | timeNormalizationAux | timeNormalizationAux | timeNormalizationAux | | | || calculated on the CPU... |
245 // | nBadHVCells || | | | | significanceAux | significanceAux | significanceAux | significanceAux | | | || calculated on the CPU... |
246 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
247 // | PTD || | | mX | mX | mX -> PTD | PTD | PTD | PTD | PTD | PTD | PTD || |
248 // | mass || | | mXAux | mXAux | mXAux -> mass | mass | mass | mass | mass | mass | mass || |
249 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
250 // | EMProbability || | | centerXAux | centerXAux | centerXAux -> matrix00Aux | matrix00Aux | matrix00Aux -> showerAxisX | showerAxisX | showerAxisX | showerAxisX | showerAxisX || |
251 // | hadWeight || | | centerYAux | centerYAux | centerYAux -> matrix11Aux | matrix11Aux | matrix11Aux -> showerAxisY | showerAxisY | showerAxisY | showerAxisY | showerAxisY || |
252 // | OOCweight || | | centerZAux | centerZAux | centerZAux -> matrix22Aux | matrix22Aux | matrix22Aux -> showerAxisZ | showerAxisZ | showerAxisZ | showerAxisZ | showerAxisZ || |
253 // | DMweight || | | clusterEnergyAux | clusterEnergyAux | clusterEnergyAux | | | lateralNormalization | lateralNormalization | lateralNormalization | || |
254 // | tileConfidenceLevel || | | clusterPhiAux | clusterPhiAux | clusterPhiAux -> matrix22 | matrix22 | matrix22 | lateralNormalizationAux | lateralNormalizationAux | lateralNormalizationAux | || |
255 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
256 // | secondTime || | | | | secondTime | secondTime | secondTime | secondTime | secondTime | secondTime | secondTime || |
257 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
258 // | nCellSampling [00] || | | nCellSampling [00] | nCellSampling [00] | nCellSampling [00] | nCellSampling [00] | nCellSampling [00] | nCellSampling [00] | nCellSampling [00] | nCellSampling [00] | nCellSampling [00] || |
259 // | nCellSampling [01] || | | nCellSampling [01] | nCellSampling [01] | nCellSampling [01] | nCellSampling [01] | nCellSampling [01] | nCellSampling [01] | nCellSampling [01] | nCellSampling [01] | nCellSampling [01] || |
260 // | nCellSampling [02] || | | nCellSampling [02] | nCellSampling [02] | nCellSampling [02] | nCellSampling [02] | nCellSampling [02] | nCellSampling [02] | nCellSampling [02] | nCellSampling [02] | nCellSampling [02] || |
261 // | nCellSampling [03] || | | nCellSampling [03] | nCellSampling [03] | nCellSampling [03] | nCellSampling [03] | nCellSampling [03] | nCellSampling [03] | nCellSampling [03] | nCellSampling [03] | nCellSampling [03] || |
262 // | nCellSampling [04] || | | nCellSampling [04] | nCellSampling [04] | nCellSampling [04] | nCellSampling [04] | nCellSampling [04] | nCellSampling [04] | nCellSampling [04] | nCellSampling [04] | nCellSampling [04] || |
263 // | nCellSampling [05] || | | nCellSampling [05] | nCellSampling [05] | nCellSampling [05] | nCellSampling [05] | nCellSampling [05] | nCellSampling [05] | nCellSampling [05] | nCellSampling [05] | nCellSampling [05] || |
264 // | nCellSampling [06] || | | nCellSampling [06] | nCellSampling [06] | nCellSampling [06] | nCellSampling [06] | nCellSampling [06] | nCellSampling [06] | nCellSampling [06] | nCellSampling [06] | nCellSampling [06] || |
265 // | nCellSampling [07] || | | nCellSampling [07] | nCellSampling [07] | nCellSampling [07] | nCellSampling [07] | nCellSampling [07] | nCellSampling [07] | nCellSampling [07] | nCellSampling [07] | nCellSampling [07] || |
266 // | nCellSampling [08] || | | nCellSampling [08] | nCellSampling [08] | nCellSampling [08] | nCellSampling [08] | nCellSampling [08] | nCellSampling [08] | nCellSampling [08] | nCellSampling [08] | nCellSampling [08] || |
267 // | nCellSampling [09] || | | nCellSampling [09] | nCellSampling [09] | nCellSampling [09] | nCellSampling [09] | nCellSampling [09] | nCellSampling [09] | nCellSampling [09] | nCellSampling [09] | nCellSampling [09] || |
268 // | nCellSampling [10] || | | nCellSampling [10] | nCellSampling [10] | nCellSampling [10] | nCellSampling [10] | nCellSampling [10] | nCellSampling [10] | nCellSampling [10] | nCellSampling [10] | nCellSampling [10] || |
269 // | nCellSampling [11] || | | nCellSampling [11] | nCellSampling [11] | nCellSampling [11] | nCellSampling [11] | nCellSampling [11] | nCellSampling [11] | nCellSampling [11] | nCellSampling [11] | nCellSampling [11] || |
270 // | nCellSampling [12] || | | nCellSampling [12] | nCellSampling [12] | nCellSampling [12] | nCellSampling [12] | nCellSampling [12] | nCellSampling [12] | nCellSampling [12] | nCellSampling [12] | nCellSampling [12] || |
271 // | nCellSampling [13] || | | nCellSampling [13] | nCellSampling [13] | nCellSampling [13] | nCellSampling [13] | nCellSampling [13] | nCellSampling [13] | nCellSampling [13] | nCellSampling [13] | nCellSampling [13] || |
272 // | nCellSampling [14] || | | nCellSampling [14] | nCellSampling [14] | nCellSampling [14] | nCellSampling [14] | nCellSampling [14] | nCellSampling [14] | nCellSampling [14] | nCellSampling [14] | nCellSampling [14] || |
273 // | nCellSampling [15] || | | nCellSampling [15] | nCellSampling [15] | nCellSampling [15] | nCellSampling [15] | nCellSampling [15] | nCellSampling [15] | nCellSampling [15] | nCellSampling [15] | nCellSampling [15] || |
274 // | nCellSampling [16] || | | nCellSampling [16] | nCellSampling [16] | nCellSampling [16] | nCellSampling [16] | nCellSampling [16] | nCellSampling [16] | nCellSampling [16] | nCellSampling [16] | nCellSampling [16] || |
275 // | nCellSampling [17] || | | nCellSampling [17] | nCellSampling [17] | nCellSampling [17] | nCellSampling [17] | nCellSampling [17] | nCellSampling [17] | nCellSampling [17] | nCellSampling [17] | nCellSampling [17] || |
276 // | nCellSampling [18] || | | nCellSampling [18] | nCellSampling [18] | nCellSampling [18] | nCellSampling [18] | nCellSampling [18] | nCellSampling [18] | nCellSampling [18] | nCellSampling [18] | nCellSampling [18] || |
277 // | nCellSampling [19] || | | nCellSampling [19] | nCellSampling [19] | nCellSampling [19] | nCellSampling [19] | nCellSampling [19] | nCellSampling [19] | nCellSampling [19] | nCellSampling [19] | nCellSampling [19] || |
278 // | nCellSampling [20] || | | nCellSampling [20] | nCellSampling [20] | nCellSampling [20] | nCellSampling [20] | nCellSampling [20] | nCellSampling [20] | nCellSampling [20] | nCellSampling [20] | nCellSampling [20] || |
279 // | nCellSampling [21] || | | nCellSampling [21] | nCellSampling [21] | nCellSampling [21] | nCellSampling [21] | nCellSampling [21] | nCellSampling [21] | nCellSampling [21] | nCellSampling [21] | nCellSampling [21] || |
280 // | nCellSampling [22] || | | nCellSampling [22] | nCellSampling [22] | nCellSampling [22] | nCellSampling [22] | nCellSampling [22] | nCellSampling [22] | nCellSampling [22] | nCellSampling [22] | nCellSampling [22] || |
281 // | nCellSampling [23] || | | nCellSampling [23] | nCellSampling [23] | nCellSampling [23] | nCellSampling [23] | nCellSampling [23] | nCellSampling [23] | nCellSampling [23] | nCellSampling [23] | nCellSampling [23] || |
282 // | nCellSampling [24] || | | nCellSampling [24] | nCellSampling [24] | nCellSampling [24] | nCellSampling [24] | nCellSampling [24] | nCellSampling [24] | nCellSampling [24] | nCellSampling [24] | nCellSampling [24] || |
283 // | nCellSampling [25] || | | nCellSampling [25] | nCellSampling [25] | nCellSampling [25] | nCellSampling [25] | nCellSampling [25] | nCellSampling [25] | nCellSampling [25] | nCellSampling [25] | nCellSampling [25] || |
284 // | nCellSampling [26] || | | nCellSampling [26] | nCellSampling [26] | nCellSampling [26] | nCellSampling [26] | nCellSampling [26] | nCellSampling [26] | nCellSampling [26] | nCellSampling [26] | nCellSampling [26] || |
285 // | nCellSampling [27] || | | nCellSampling [27] | nCellSampling [27] | nCellSampling [27] | nCellSampling [27] | nCellSampling [27] | nCellSampling [27] | nCellSampling [27] | nCellSampling [27] | nCellSampling [27] || |
286 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
287 // | nExtraCellSampling || | | sumAbsEnergyNonMomentsAux | sumAbsEnergyNonMomentsAux | sumAbsEnergyNonMomentsAux | | | nExtraCellSampling | nExtraCellSampling | nExtraCellSampling | nExtraCellSampling || |
288 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
289 // | vertexFraction || | | | | engBadCellsAux | engBadCellsAux | engBadCellsAux | engBadCellsAux | | | || |
290 // | nVertexFraction || | | | | badLArQFracAux | badLArQFracAux | badLArQFracAux | badLArQFracAux | | | || |
291 // | etaCaloFrame || | | energyDensityNormalization | energyDensityNormalization |enrgyDnstyNorm > avgLArQAux | avgLArQAux | avgLArQAux | avgLArQAux | | | || |
292 // | phiCaloFrame || | | | | timeAux | timeAux | timeAux | timeAux | | | || |
293 // | eta1CaloFrame || | | | | secondTimeAux | secondTimeAux | secondTimeAux | secondTimeAux | | | || |
294 // | phi1CaloFrame || | | | | averageTileQNorm | averageTileQNorm | averageTileQNorm | averageTileQNorm | | | || |
295 // | eta2CaloFrame || | | | | averageTileQNormAux | averageTileQNormAux | averageTileQNormAux | averageTileQNormAux | | | || |
296 // | phi2CaloFrame || | | | | timeNormalization | timeNormalization | timeNormalization | timeNormalization | | | || |
297 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
298 // | engCalibTot || <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> || <struct's extra cell info> |
299 // | engCalibOutL || <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> || <struct's extra cell info> |
300 // | engCalibOutM || <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> || <struct's extra cell info> |
301 // | engCalibOutT || <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> || <struct's extra cell info> |
302 // | engCalibDeadL || <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> || <struct's extra cell info> |
303 // | engCalibDeadM || <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> | <struct's extra cell info> || <struct's extra cell info> |
304 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
305 // | engCalibDeadT || | | | | matrix10 | matrix10 | matrix10 | | | | || |
306 // | engCalibEMB0 || | | | | matrix20 | matrix20 | matrix20 | lateralAux | lateralAux | lateralAux | || |
307 // | engCalibEME0 || | | | | matrix21 | matrix21 | matrix21 | longitudinalAux | longitudinalAux | longitudinalAux | || |
308 // | engCalibTileG3 || | | firstEngDensAux | firstEngDensAux |firstEngDensAux>avgLArQNorm | averageLArQNorm | averageLArQNorm |avgLArQNorm>secondLambdaAux | secondLambdaAux | secondLambdaAux | || |
309 // | engCalibDeadTot || | | | | sumSquareEnergies | sumSquareEnergies | sumSquareEnergies | secondRAux | secondRAux | secondRAux | || |
310 // | engCalibDeadEMB0 || | | | | sumSquareEnergiesAux | sumSquareEnergiesAux | sumSquareEnergiesAux | | | | || |
311 // | engCalibDeadTile0 || | | maxCellEnergyAndCell_2 | maxCellEnergyAndCell_2 |mCEAC2>mxSgnificanceAndSmpl2|maxSignificanceAndSampling_2|maxSignificanceAndSampling_2|maxSignificanceAndSampling_2| | | || |
312 // | engCalibDeadTileG3 || | | maxCellEnergyAndCell_1 | maxCellEnergyAndCell_1 |mCEAC1>mxSgnificanceAndSmpl1|maxSignificanceAndSampling_1|maxSignificanceAndSampling_1|maxSignificanceAndSampling_1| | | || |
313 // | engCalibDeadEME0 || | |secondMaxCellEnergyAndCell_2|secondMaxCellEnergyAndCell_2|sMCEAC2>maxAndScondMaxCells2| maxAndSecondMaxCells_2 | maxAndSecondMaxCells_2 | maxAndSecondMaxCells_2 | maxAndSecondMaxCells_2 | | || |
314 // | engCalibDeadHEC0 || | |secondMaxCellEnergyAndCell_1|secondMaxCellEnergyAndCell_1|sMCEAC2>maxAndScondMaxCells1| maxAndSecondMaxCells_1 | maxAndSecondMaxCells_1 | maxAndSecondMaxCells_1 | maxAndSecondMaxCells_1 | | || |
315 // | engCalibDeadFCAL || | | clusterEtaAux | clusterEtaAux | clusterEtaAux | | | | | | || |
316 // | engCalibDeadLeakage || | | mY | mY | mY -> PTDAux | PTDAux | PTDAux | PTDAux | | | || |
317 // | engCalibDeadUnclass || | | mZ | mZ | mZ > numPositiveEnergyCells| numPositiveEnergyCells | numPositiveEnergyCells | | | | || |
318 // | engCalibFracEM || | | | | avgTileQAux | avgTileQAux | avgTileQAux | avgTileQAux | | | || |
319 // | engCalibFracHad || | | sumAbsEnergyNonMoments | sumAbsEnergyNonMoments | sumAbsEnergyNonMoments | | | longitudinalNormalization | longitudinalNormalization | longitudinalNormalization | || |
320 // | engCalibFracRest || | | seedCellPhi | seedCellPhi | | | |longitudinalNormalizationAux|longitudinalNormalizationAux|longitudinalNormalizationAux| || |
321 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
322 // | CLUSTER STORAGE VARIABLE || Isolation Cluster Pass | Isolation Cell Pass | Zeroth Cluster Pass | First Cell Pass | First Cluster Pass | Second Cell Pass | Shower Axis Pass | Second Cluster Pass | Third Cell Pass | Third Cluster Pass | Final Cluster Pass || Comments |
323 // +--------------------------++----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------+----------------------------++----------------------------+
324 //
325
326
327 CALORECGPU_TEMP2DARR_1 ( numberEmptySamplings, etaPerSample, int);
328 CALORECGPU_TEMP2DARR_1 ( numberNonEmptySamplings, maxPhiPerSample, int);
329 CALORECGPU_TEMP2DARR_1 ( maxMomentsEnergyPerSample, maxEtaPerSample, unsigned int);
330
331
332 CALORECGPU_TEMP2DARR_1 ( absoluteEnergyPerSample, maxEPerSample, float);
333 CALORECGPU_TEMP2DARR_1 ( absoluteEnergyPerSampleAux, energyPerSample, float);
334 CALORECGPU_TEMP2DARR_1 ( phiPerSampleAux, maxPhiPerSample, float);
335 CALORECGPU_TEMP2DARR_1 ( etaPerSampleAux, maxEtaPerSample, float);
336 CALORECGPU_TEMPARR_1 ( seedCellPhi, engCalibFracRest, float);
337 CALORECGPU_TEMPARR_1 ( clusterEnergyAux, DMweight, float);
338 CALORECGPU_TEMPARR_1 ( clusterPhiAux, tileConfidenceLevel, float);
339 CALORECGPU_TEMPARR_1 ( clusterEtaAux, engCalibDeadFCAL, float);
340 CALORECGPU_TEMPARR_1 ( centerXAux, EMProbability, float);
341 CALORECGPU_TEMPARR_1 ( centerYAux, hadWeight, float);
342 CALORECGPU_TEMPARR_1 ( centerZAux, OOCweight, float);
343 CALORECGPU_TEMPARR_1 ( firstPhiAux, secondR, float);
344 CALORECGPU_TEMPARR_1 ( firstEtaAux, longitudinal, float);
345 CALORECGPU_TEMPARR_1 ( engPosAux, engBadHVCells, float);
346 CALORECGPU_TEMPARR_1 ( engFracEMAux, centerMag, float);
347 CALORECGPU_TEMPARR_1 ( firstEngDensAux, engCalibTileG3, float);
348 CALORECGPU_TEMPARR_1 ( secondEngDensAux, lateral, float);
349 CALORECGPU_TEMPARR_1 ( energyDensityNormalization, etaCaloFrame, float);
350 CALORECGPU_TEMPARR_1 ( energyDensityNormalizationAux, secondLambda, float);
351 CALORECGPU_TEMPARR_1 ( sumAbsEnergyNonMoments, engCalibFracHad, float);
352 CALORECGPU_TEMPARR_1 ( sumAbsEnergyNonMomentsAux, nExtraCellSampling, float);
353 CALORECGPU_TEMPARR_1 ( mX, PTD, float);
354 CALORECGPU_TEMPARR_1 ( mXAux, mass, float);
355 CALORECGPU_TEMPARR_1 ( mY, engCalibDeadLeakage, float);
356 CALORECGPU_TEMPARR_1 ( mYAux, cellSignificance, float);
357 CALORECGPU_TEMPARR_1 ( mZ, engCalibDeadUnclass, float);
358 CALORECGPU_TEMPARR_1 ( mZAux, cellSigSampling, float);
359 CALORECGPU_TEMPARR_2 ( maxCellEnergyAndCell, engCalibDeadTileG3, engCalibDeadTile0, unsigned long long);
360 CALORECGPU_TEMPARR_2 ( secondMaxCellEnergyAndCell, engCalibDeadHEC0, engCalibDeadEME0, unsigned long long);
361 CALORECGPU_TEMPARR_2 ( maxAndSecondMaxCells, engCalibDeadHEC0, engCalibDeadEME0, unsigned long long);
362
363 CALORECGPU_TEMPARR_1 ( timeAux, phiCaloFrame, float);
364 CALORECGPU_TEMPARR_1 ( secondTimeAux, eta1CaloFrame, float);
365 CALORECGPU_TEMPARR_1 ( significanceAux, nBadHVCells, float);
366 CALORECGPU_TEMPARR_1 ( PTDAux, engCalibDeadLeakage, float);
367 CALORECGPU_TEMPARR_1 ( engBadCellsAux, vertexFraction, float);
368 CALORECGPU_TEMPARR_1 ( badCellsCorrEAux, centerLambda, float);
369 CALORECGPU_TEMPARR_1 ( badLArQFracAux, nVertexFraction, float);
370 CALORECGPU_TEMPARR_1 ( avgLArQAux, etaCaloFrame, float);
371 CALORECGPU_TEMPARR_1 ( avgTileQAux, engCalibFracEM, float);
372 CALORECGPU_TEMPARR_1 ( numPositiveEnergyCells, engCalibDeadUnclass, int);
373 CALORECGPU_TEMPARR_1 ( sumSquareEnergies, engCalibDeadTot, float);
374 CALORECGPU_TEMPARR_1 ( sumSquareEnergiesAux, engCalibDeadEMB0, float);
375 CALORECGPU_TEMPARR_1 ( matrix00, longitudinal, float);
376 CALORECGPU_TEMPARR_1 ( matrix10, engCalibDeadT, float);
377 CALORECGPU_TEMPARR_1 ( matrix20, engCalibEMB0, float);
378 CALORECGPU_TEMPARR_1 ( matrix11, secondR, float);
379 CALORECGPU_TEMPARR_1 ( matrix21, engCalibEME0, float);
380 CALORECGPU_TEMPARR_1 ( matrix22, tileConfidenceLevel, float);
381 CALORECGPU_TEMPARR_1 ( matrix00Aux, EMProbability, float);
382 CALORECGPU_TEMPARR_1 ( matrix10Aux, deltaPhi, float);
383 CALORECGPU_TEMPARR_1 ( matrix20Aux, deltaTheta, float);
384 CALORECGPU_TEMPARR_1 ( matrix11Aux, hadWeight, float);
385 CALORECGPU_TEMPARR_1 ( matrix21Aux, deltaAlpha, float);
386 CALORECGPU_TEMPARR_1 ( matrix22Aux, OOCweight, float);
387 CALORECGPU_TEMPARR_1 ( timeNormalization, phi2CaloFrame, float);
388 CALORECGPU_TEMPARR_1 ( timeNormalizationAux, engBadHVCells, float);
389 CALORECGPU_TEMPARR_1 ( averageLArQNorm, engCalibTileG3, float);
390 CALORECGPU_TEMPARR_1 ( averageLArQNormAux, secondLambda, float);
391 CALORECGPU_TEMPARR_1 ( averageTileQNorm, phi1CaloFrame, float);
392 CALORECGPU_TEMPARR_1 ( averageTileQNormAux, eta2CaloFrame, float);
393 CALORECGPU_TEMPARR_2 ( maxSignificanceAndSampling, engCalibDeadTileG3, engCalibDeadTile0, unsigned long long);
394 CALORECGPU_TEMPARR_1 ( showerAxisX, EMProbability, float);
395 CALORECGPU_TEMPARR_1 ( showerAxisY, hadWeight, float);
396 CALORECGPU_TEMPARR_1 ( showerAxisZ, OOCweight, float);
397
398 CALORECGPU_TEMP2DARR_2 ( maxEnergyAndCellPerSample, maxEtaPerSample, maxPhiPerSample, unsigned long long);
399 CALORECGPU_TEMP2DARR_1 ( energyPerSampleAux, maxEPerSample, float);
400 CALORECGPU_TEMPARR_1 ( lateralAux, engCalibEMB0, float);
401 CALORECGPU_TEMPARR_1 ( longitudinalAux, engCalibEME0, float);
402 CALORECGPU_TEMPARR_1 ( secondLambdaAux, engCalibTileG3, float);
403 CALORECGPU_TEMPARR_1 ( secondRAux, engCalibDeadTot, float);
404 CALORECGPU_TEMPARR_1 ( lateralNormalization, DMweight, float);
405 CALORECGPU_TEMPARR_1 ( lateralNormalizationAux, tileConfidenceLevel, float);
406 CALORECGPU_TEMPARR_1 ( longitudinalNormalization, engCalibFracHad, float);
407 CALORECGPU_TEMPARR_1 ( longitudinalNormalizationAux, engCalibFracRest, float);
408
409 CALORECGPU_TEMP2DARR_1 ( maxECellPerSample, maxEPerSample, int);
410
411 }
412
413 template <class ... Ts>
414 struct TypeList
415 {
416 static constexpr unsigned int number = static_cast<unsigned int>(sizeof...(Ts));
417 };
418
419 template <class ... Ts, class T>
420 constexpr __host__ __device__ auto operator| (const TypeList<Ts...> & list, const TypeList<T> &)
421 {
422 if constexpr ((false || ... || std::is_same_v<T, Ts>))
423 {
424 return list;
425 }
426 else
427 {
428 return TypeList<Ts..., T> {};
429 }
430 }
431
432 template <class ... As, class ... Bs>
433 constexpr __host__ __device__ auto operator | (const TypeList<As...> & list, const TypeList<Bs...> &)
434 {
435 return (list | ... | TypeList<Bs> {});
436 }
437
438 template <class ... As, class ... Bs>
439 constexpr __host__ __device__ auto operator + (const TypeList<As...> & list, const TypeList<Bs...> &)
440 {
441 return TypeList<As..., Bs...> {};
442 }
443
444 template <class T>
445 struct is_type_list
446 {
447 static constexpr bool value = false;
448 };
449
450 template <class ... Ts>
451 struct is_type_list<TypeList<Ts...>>
452 {
453 static constexpr bool value = true;
454 };
455
456 template <class T>
457 inline constexpr bool is_type_list_v = is_type_list<T>::value;
458
459 struct Parameters
460 {
461 bool assume_complete_cells;
462 int moments_index;
463 CaloRecGPU::Helpers::CUDA_kernel_object<CaloRecGPU::ClusterInfoArr> clusters_arr;
464 const CaloRecGPU::Helpers::CUDA_kernel_object<CaloRecGPU::CellInfoArr> cell_info_arr;
465 const CaloRecGPU::Helpers::CUDA_kernel_object<CaloRecGPU::GeometryArr> geometry;
466 const CaloRecGPU::Helpers::CUDA_kernel_object<CaloRecGPU::CellNoiseArr> noise_arr;
467 const CaloRecGPU::Helpers::CUDA_kernel_object<ClusterMomentsCalculator::ClusterMomentCalculationOptions> opts;
468 };
469
470 template <class T>
471 struct OnePassLoadHelper
472 {
473 static_assert(is_type_list_v<T>);
474
475 using ToLoad = T;
476 };
477
478 //Add property if unique
479 template <class ... Ts, class T>
480 constexpr __host__ __device__ auto operator + (const OnePassLoadHelper<TypeList<Ts...>> &, const TypeList<T> &)
481 {
482 constexpr auto combined = TypeList<Ts...> {} | TypeList<T> {};
483
484 return OnePassLoadHelper<std::decay_t<decltype(combined)>> {};
485 }
486
487 //Add any dependencies and then the desired property
488 template <class T, class NewLoad>
489 constexpr __host__ __device__ auto operator | (const OnePassLoadHelper<T> & c, const TypeList<NewLoad> & n)
490 {
491 auto combine = c | typename NewLoad::AssumedList{};
492
493 return combine + n;
494 }
495
496 //Add all properties with potential dependencies
497 template <class T, class ... Ts>
498 constexpr __host__ __device__ auto operator | (const OnePassLoadHelper<T> & c, const TypeList<Ts...> &)
499 {
500 return (c | ... | TypeList<Ts> {});
501 }
502
503 template <class T>
504 using ClusterPassLoadHelper = OnePassLoadHelper<T>;
505
506 template <class T1, class T2>
507 struct CellAndClusterPassLoadHelper
508 {
509 static_assert(is_type_list_v<T1> && is_type_list_v<T2>);
510
511 using ToLoadCells = T1;
512 using ToLoadClusters = T2;
513 };
514
515 //Use | to add cell loading with dependencies
516 //Use ^ to add cluster loading with dependencies
517
518 template <class T1, class T2, class NewCellLoad>
519 constexpr __host__ __device__ auto operator | (const CellAndClusterPassLoadHelper<T1, T2> & c, const TypeList<NewCellLoad> & n)
520 {
521 auto combine_cell = OnePassLoadHelper<T1> {} | n;
522
523 return CellAndClusterPassLoadHelper<typename std::decay_t<decltype(combine_cell)>::ToLoad, T2> {};
524 }
525
526 template <class T1, class T2, class ... Ts>
527 constexpr __host__ __device__ auto operator | (const CellAndClusterPassLoadHelper<T1, T2> & c, const TypeList<Ts...> &)
528 {
529 return (c | ... | TypeList<Ts> {});
530 }
531
532 template <class T1, class T2, class NewClusterLoad>
533 constexpr __host__ __device__ auto operator ^ (const CellAndClusterPassLoadHelper<T1, T2> & c, const TypeList<NewClusterLoad> & n)
534 {
535 auto combine_dependencies = (c | typename NewClusterLoad::AssumedPreviousList{}) ^ typename NewClusterLoad::AssumedList{};
536
537 using DepType = std::decay_t<decltype(combine_dependencies)>;
538
539 auto combine_clusters = typename DepType::ToLoadClusters{} | n;
540
541 return CellAndClusterPassLoadHelper<typename DepType::ToLoadCells, std::decay_t<decltype(combine_clusters)>> {};
542 }
543
544 template <class T1, class T2, class ... Ts>
545 constexpr __host__ __device__ auto operator ^ (const CellAndClusterPassLoadHelper<T1, T2> & c, const TypeList<Ts...> &)
546 {
547 return (c ^ ... ^ TypeList<Ts> {});
548 }
549
550 template <class T>
551 struct ClusterLoader;
552
553 template <class ... Ts>
554 struct ClusterLoader<ClusterPassLoadHelper<TypeList<Ts...>>> : Ts...
555 {
556 __host__ __device__ ClusterLoader(Parameters p, const int idx):
557 Ts(*this, p, idx)...
558 {
559 }
560 };
561
562 template <class T>
563 struct CellClusterLoader;
564
565 struct WeightCarrier
566 {
567 float weight;
568
569 __host__ __device__ WeightCarrier(const float w): weight(w)
570 {
571 }
572 };
573
574 template <class ... CellLoad, class ... ClusterLoad>
575 struct CellClusterLoader<CellAndClusterPassLoadHelper<TypeList<CellLoad...>, TypeList<ClusterLoad...>>>: WeightCarrier, CellLoad..., ClusterLoad...
576 {
577 __host__ __device__ CellClusterLoader(Parameters p, const int cell, const int cluster, const float weight):
578 WeightCarrier(weight), CellLoad(*this, p, cell)..., ClusterLoad(*this, p, cluster)...
579 {
580 }
581 };
582
583 template <class ... Moments>
584 __host__ __device__ void one_cluster_before_pass(const TypeList<Moments...> &, const int cluster, Parameters p)
585 {
586 constexpr auto to_load = (ClusterPassLoadHelper<TypeList<>> {} | ... | typename Moments::BeforeLoading{});
587
588 ClusterLoader<std::decay_t<decltype(to_load)>> data{p, cluster};
589
590 (Moments::before(p, data, cluster), ...);
591 }
592
593 template <class ... Moments>
594 __host__ __device__ void one_cluster_after_pass(const TypeList<Moments...> &, const int cluster, Parameters p)
595 {
596 constexpr auto to_load = (ClusterPassLoadHelper<TypeList<>> {} | ... | typename Moments::AfterLoading{});
597
598 ClusterLoader<std::decay_t<decltype(to_load)>> data{p, cluster};
599
600 (Moments::after(p, data, cluster), ...);
601 }
602
603 template <class ... Moments>
604 __host__ __device__ void one_cell_pass(const TypeList<Moments ...> &, const int cell, const int cluster, const float weight, Parameters p)
605 {
606 constexpr auto to_load = ((CellAndClusterPassLoadHelper<TypeList<>, TypeList<>> {} | ... | typename Moments::CellLoading{}) ^ ... ^ typename Moments::ClusterLoading{});
607
608 CellClusterLoader<std::decay_t<decltype(to_load)>> data{p, cell, cluster, weight};
609
610 (Moments::per_cell(p, data, cell, cluster), ...);
611
612 }
613
617 template <class ... ThisStepMomentLists, class ... NextStepMomentLists>
618 __host__ __device__ void do_cluster_pass(const TypeList<ThisStepMomentLists...> &, const TypeList<NextStepMomentLists...> &, const int cluster, Parameters p)
619 {
620 auto helper_1 = [&](const auto & list, const int count)
621 {
622 if (p.moments_index == count)
623 {
624 one_cluster_after_pass(list, cluster, p);
625 }
626 };
627 auto helper_2 = [&](const auto & list, const int count)
628 {
629 if (p.moments_index == count)
630 {
631 one_cluster_before_pass(list, cluster, p);
632 }
633 };
634
635 int count = 0;
636
637 (helper_1(ThisStepMomentLists{}, count++), ...);
638
639 count = 0;
640
641 (helper_2(NextStepMomentLists{}, count++), ...);
642 }
643
644 template <class... MomentLists>
645 __host__ __device__ void do_cell_pass(const TypeList<MomentLists...> &, const int cell, const int cluster, const float weight, Parameters p)
646 {
647 auto helper = [&](const auto & list, const int count)
648 {
649 if (p.moments_index == count)
650 {
651 one_cell_pass(list, cell, cluster, weight, p);
652 }
653 };
654
655 int count = 0;
656
657 (helper(MomentLists{}, count++), ...);
658 }
659
660
661 //-------------------------------
662
663 namespace ToLoad
664 {
665
666#define CALORECGPU_CMC_EXPAND(...) __VA_ARGS__
667
668
669#define CALORECGPU_CMC_LOAD(NAME, NEEDED, PREVNEEDED, VARS, INIT) \
670 struct NAME \
671 { \
672 using AssumedList = TypeList<CALORECGPU_CMC_EXPAND NEEDED >; \
673 using AssumedPreviousList = TypeList<CALORECGPU_CMC_EXPAND PREVNEEDED >; \
674 CALORECGPU_CMC_EXPAND VARS \
675 template <class Final> __device__ NAME(const Final & f, Parameters p, const int idx) { CALORECGPU_CMC_EXPAND INIT } \
676 }
677 //Of course, any circular dependencies will lead to infinite loops during compilation!
678
679
680#define CALORECGPU_CMC_LOAD_SIMPLE_CELL_INFO(NAME, VARNAME, PROPNAME) \
681 CALORECGPU_CMC_LOAD(NAME, \
682 (), \
683 (), \
684 (std::decay_t<decltype(std::declval<CaloRecGPU::CellInfoArr>().PROPNAME[0])> VARNAME;), \
685 (VARNAME = p.cell_info_arr->PROPNAME[idx];) \
686 );
687
688#define CALORECGPU_CMC_LOAD_SIMPLE_GEOMETRY_INFO(NAME, VARNAME, PROPNAME) \
689 CALORECGPU_CMC_LOAD(NAME, \
690 (CellHashID), \
691 (), \
692 (std::decay_t<decltype(std::declval<CaloRecGPU::GeometryArr>().PROPNAME[0])> VARNAME;), \
693 (VARNAME = p.geometry->PROPNAME[f.hash_ID];) \
694 );
695
696#define CALORECGPU_CMC_LOAD_SIMPLE_CLUSTER_INFO(NAME, VARNAME, PROPNAME) \
697 CALORECGPU_CMC_LOAD(NAME, \
698 (), \
699 (), \
700 (std::decay_t<decltype(std::declval<CaloRecGPU::ClusterInfoArr>().PROPNAME[0])> VARNAME;), \
701 (VARNAME = p.clusters_arr->PROPNAME[idx];) \
702 );
703
704#define CALORECGPU_CMC_LOAD_SIMPLE_MOMENT_INFO(NAME, VARNAME, PROPNAME) \
705 CALORECGPU_CMC_LOAD(NAME, \
706 (), \
707 (), \
708 (std::decay_t<decltype(std::declval<CaloRecGPU::ClusterInfoArr>().moments.PROPNAME[0])> VARNAME;), \
709 (VARNAME = p.clusters_arr->moments.PROPNAME[idx];) \
710 );
711
712 //Warning! These require either a CellSampling or a SamplingFromMomentIndex
713 //to be calculated!
714#define CALORECGPU_CMC_LOAD_SIMPLE_PER_SAMPLING_MOMENT_INFO(NAME, VARNAME, PROPNAME) \
715 CALORECGPU_CMC_LOAD(NAME, \
716 (), \
717 (), \
718 (std::decay_t<decltype(std::declval<CaloRecGPU::ClusterInfoArr>().moments.PROPNAME[0][0])> VARNAME;), \
719 (VARNAME = p.clusters_arr->moments.PROPNAME[f.sampling][idx];) \
720 );
721
722#define CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(NAME, VARNAME, PROPNAME) \
723 CALORECGPU_CMC_LOAD(NAME, \
724 (), \
725 (), \
726 (std::decay_t<decltype(CMCTemporaries::PROPNAME(std::declval<CaloRecGPU::ClusterInfoArr *>(),0))> VARNAME;), \
727 (VARNAME = CMCTemporaries::PROPNAME(p.clusters_arr, idx);) \
728 );
729
730 //Warning! These require either a CellSampling or a SamplingFromMomentIndex
731 //to be calculated!
732#define CALORECGPU_CMC_LOAD_SIMPLE_PER_SAMPLING_TEMPORARY_INFO(NAME, VARNAME, PROPNAME) \
733 CALORECGPU_CMC_LOAD(NAME, \
734 (), \
735 (), \
736 (std::decay_t<decltype(CMCTemporaries::PROPNAME(std::declval<CaloRecGPU::ClusterInfoArr *>(),0,0))> VARNAME;), \
737 (VARNAME = CMCTemporaries::PROPNAME(p.clusters_arr, f.sampling, idx);) \
738 );
739
740 //+---------------------------------------------------+
741 //| Sampling Variables: |
742 //+---------------------------------------------------+
743
744 CALORECGPU_CMC_LOAD
745 (SamplingFromMomentIndex,
746 (),
747 (),
748 (int sampling;),
749 (sampling = p.moments_index;)
750 );
751
752 CALORECGPU_CMC_LOAD
753 (CellHashID,
754 (),
755 (),
756 (int hash_ID;),
757 (hash_ID = p.cell_info_arr->get_hash_ID(idx, p.assume_complete_cells);)
758 );
759
760 CALORECGPU_CMC_LOAD
761 (CellSampling,
762 (CellHashID),
763 (),
764 (int sampling;),
765 (sampling = p.geometry->sampling(f.hash_ID);)
766 );
767
768 //+---------------------------------------------------+
769 //| Cell Only Info: |
770 //+---------------------------------------------------+
771
772 CALORECGPU_CMC_LOAD_SIMPLE_GEOMETRY_INFO(CellX, x, x);
773 CALORECGPU_CMC_LOAD_SIMPLE_GEOMETRY_INFO(CellY, y, y);
774 CALORECGPU_CMC_LOAD_SIMPLE_GEOMETRY_INFO(CellZ, z, z);
775 CALORECGPU_CMC_LOAD_SIMPLE_GEOMETRY_INFO(CellEta, eta, eta);
776 CALORECGPU_CMC_LOAD_SIMPLE_GEOMETRY_INFO(CellPhi, phi, phi);
777 CALORECGPU_CMC_LOAD_SIMPLE_GEOMETRY_INFO(CellVolume, volume, volume);
778
779 CALORECGPU_CMC_LOAD_SIMPLE_CELL_INFO(CellEnergy, energy, energy);
780
781 CALORECGPU_CMC_LOAD
782 (CellAbsEnergy,
783 (CellEnergy),
784 (),
785 (float abs_energy;),
786 (abs_energy = fabsf(f.energy);)
787 );
788
789 CALORECGPU_CMC_LOAD
790 (CellMomentsEnergy,
791 (CellEnergy, CellAbsEnergy),
792 (),
793 (float moments_energy;),
794 (moments_energy = (p.opts->use_abs_energy || f.energy > 0.f) ? f.abs_energy : 0.f;)
795 );
796
797 CALORECGPU_CMC_LOAD_SIMPLE_CELL_INFO(CellTime, time, time);
798 CALORECGPU_CMC_LOAD_SIMPLE_CELL_INFO(CellGain, gain, gain);
799
800 CALORECGPU_CMC_LOAD
801 (CellQualityProvenance,
802 (),
803 (),
804 (CaloRecGPU::QualityProvenance qp{0};),
805 (qp = p.cell_info_arr->qualityProvenance[idx];)
806 );
807
808 CALORECGPU_CMC_LOAD
809 (CellIsTile,
810 (CellHashID),
811 (),
812 (bool is_tile;),
813 (is_tile = p.geometry->is_tile(f.hash_ID);)
814 );
815
816 CALORECGPU_CMC_LOAD
817 (CellIsBad,
818 (CellIsTile, CellQualityProvenance),
819 (),
820 (bool is_bad;),
821 (is_bad = p.cell_info_arr->is_bad_general(f.is_tile, f.qp, false);)
822 );
823
824 CALORECGPU_CMC_LOAD
825 (CellNoise,
826 (CellHashID, CellIsTile, CellGain, CellEnergy),
827 (),
828 (float noise;),
829 (noise = ( f.is_tile && p.opts->use_two_gaussian_noise ?
830 p.noise_arr->get_double_gaussian_noise(f.hash_ID, f.gain, f.energy) :
831 p.noise_arr->get_noise(f.hash_ID, f.gain) );)
832 );
833
834 CALORECGPU_CMC_LOAD
835 (CellLArQCheck,
836 (CellIsBad, CellIsTile, CellQualityProvenance),
837 (),
838 (bool LArQ_cell_check;),
839 (LArQ_cell_check = !f.is_bad && !f.is_tile && ((f.qp.provenance() & 0x2800U) == 0x2000U);)
840 );
841
842 CALORECGPU_CMC_LOAD
843 (CellTileQCheck,
844 (CellIsBad, CellIsTile, CellQualityProvenance),
845 (),
846 (bool TileQ_cell_check;),
847 (TileQ_cell_check = !f.is_bad && f.is_tile && f.qp.tile_qual1() != 0xFFU && f.qp.tile_qual2() != 0xFFU;)
848 );
849
850
851 CALORECGPU_CMC_LOAD
852 (CellTimeMomentsCheck,
853 (CellIsTile, CellQualityProvenance, CellSampling),
854 (),
855 (bool time_moments_check;),
856 (time_moments_check = ( (f.is_tile && (f.qp.provenance() & 0x8080U)) ||
857 (!f.is_tile && (f.qp.provenance() & 0x2000U)) ) &&
858 f.sampling != CaloSampling::PreSamplerB && f.sampling != CaloSampling::PreSamplerE; )
859 );
860
861
862
863 //+---------------------------------------------------+
864 //| Cell-In-Cluster-Dependent Info: |
865 //+---------------------------------------------------+
866
867 CALORECGPU_CMC_LOAD
868 (WeightedEnergy,
869 (),
870 (CellMomentsEnergy),
871 (float weighted_energy;),
872 (weighted_energy = f.moments_energy * f.weight;)
873 );
874
875 CALORECGPU_CMC_LOAD
876 (SquareWeightedEnergy,
877 (WeightedEnergy),
878 (),
879 (float square_w_E;),
880 (square_w_E = f.weighted_energy * f.weighted_energy;)
881 );
882
883 CALORECGPU_CMC_LOAD
884 (WeightedEnergyOverVolume,
885 (WeightedEnergy),
886 (CellVolume),
887 (float w_E_over_V;),
888 (w_E_over_V = (f.volume > 0.f ? f.weighted_energy / f.volume : 1.f);)
889 );
890
891 CALORECGPU_CMC_LOAD
892 (WeightedCellPositionNormalization,
893 (WeightedEnergy),
894 (CellX, CellY, CellZ),
895 (float r_dir;
896 float w_E_r_dir;
897 ),
898 (const float r_dir_base = rnorm3df(f.x, f.y, f.z);
899 r_dir = isinf(r_dir_base) ? 0.f : r_dir_base;
900 w_E_r_dir = f.weighted_energy * f.r_dir;
901 )
902 );
903
904 CALORECGPU_CMC_LOAD
905 (WeightedEnergyOrNegative,
906 (),
907 (CellEnergy),
908 (float weighted_energy_or_negative;),
909 (weighted_energy_or_negative = (p.opts->use_abs_energy ? fabsf(f.energy) : f.energy) * f.weight;)
910 );
911
912 CALORECGPU_CMC_LOAD
913 (SquareWeightedEnergyOrNegative,
914 (WeightedEnergyOrNegative),
915 (),
916 (float square_w_E_or_neg;),
917 (square_w_E_or_neg = f.weighted_energy_or_negative * f.weighted_energy_or_negative;)
918 );
919
920 CALORECGPU_CMC_LOAD
921 (WeightedNonMomentsEnergy,
922 (),
923 (CellEnergy),
924 (float normE;),
925 (normE = f.weight * f.energy;)
926 );
927
928 CALORECGPU_CMC_LOAD
929 (SquaredWeightedNonMomentsEnergy,
930 (WeightedNonMomentsEnergy),
931 (),
932 (float squared_normE;),
933 (squared_normE = f.normE * f.normE;)
934 );
935
936 struct CenterX;
937 struct CenterY;
938 struct CenterZ;
939 struct ShowerAxisX;
940 struct ShowerAxisY;
941 struct ShowerAxisZ;
942
943 // Possible simplification if we need more performance
944 // when calculating both Lambda and R (but probably worse accuracy):
945 //
946 //d\vec{v} = \vec{r}_{cell} - \vec{r}_{center}
947 //
948 // r = ||d\vec{v} \cross \vec{axis}||
949 // lambda = d\vec{v} \dot \vec{axis}
950 //
951 // lambda = ||d\vec{v}|| cos (\theta)
952 // r = ||d\vec{v}|| |sin(\theta)|
953 //
954 // |sin(\theta)| = sqrt(1 - cos(\theta)^2)
955
956 CALORECGPU_CMC_LOAD
957 (Deltas,
958 (CenterX, CenterY, CenterZ),
959 (CellX, CellY, CellZ),
960 (float dx, dy, dz;),
961 (dx = f.x - f.center_x;
962 dy = f.y - f.center_y;
963 dz = f.z - f.center_z;
964 )
965 );
966
967 CALORECGPU_CMC_LOAD
968 (Lambda,
969 (Deltas, ShowerAxisX, ShowerAxisY, ShowerAxisZ),
970 (),
971 (float lambda;),
972 (lambda = CaloRecGPU::Helpers::corrected_dot_product(f.dx, f.dy, f.dz, f.axis_x, f.axis_y, f.axis_z);)
973 );
974
975 CALORECGPU_CMC_LOAD
976 (R,
977 (Deltas, ShowerAxisX, ShowerAxisY, ShowerAxisZ),
978 (),
979 (float r;),
980 (r = CaloRecGPU::Helpers::corrected_magn_cross_product(f.dx, f.dy, f.dz, f.axis_x, f.axis_y, f.axis_z);)
981 );
982
983 //+---------------------------------------------------+
984 //| Cluster Info: |
985 //+---------------------------------------------------+
986
987 CALORECGPU_CMC_LOAD_SIMPLE_CLUSTER_INFO(SeedCellIndex, seed_cell, seedCellIndex);
988
989 CALORECGPU_CMC_LOAD_SIMPLE_MOMENT_INFO(SumEnergies, sum_energies, engPos);
990
991 CALORECGPU_CMC_LOAD
992 (ReverseSumEnergies,
993 (SumEnergies),
994 (),
995 (float rev_sum_energies;),
996 (rev_sum_energies = 1.0f / (f.sum_energies > 0.f ? f.sum_energies : 1.f);)
997 );
998
999 CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(SeedCellPhi, phi_0, seedCellPhi);
1000
1001 CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(EnergyDensityNormalizationBase, energy_density_norm_base, energyDensityNormalization);
1002 CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(EnergyDensityNormalizationCorr, energy_density_norm_corr, energyDensityNormalizationAux);
1003
1004 CALORECGPU_CMC_LOAD
1005 (EnergyDensityNormalization,
1006 (EnergyDensityNormalizationBase, EnergyDensityNormalizationCorr),
1007 (),
1008 (float energy_density_norm;),
1009 (energy_density_norm = f.energy_density_norm_base + f.energy_density_norm_corr;)
1010 );
1011
1012 CALORECGPU_CMC_LOAD
1013 (ReverseEnergyDensityNormalization,
1014 (EnergyDensityNormalization),
1015 (),
1016 (float rev_energy_density_norm;),
1017 (rev_energy_density_norm = 1.0f / (f.energy_density_norm > 0.f ? f.energy_density_norm : 1.f);)
1018 );
1019
1020 CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(ClusterMaxCellEnergyAndCell, max_E_and_cell, maxCellEnergyAndCell);
1021
1022 CALORECGPU_CMC_LOAD
1023 (ClusterCellWithMaxEnergy,
1024 (ClusterMaxCellEnergyAndCell),
1025 (),
1026 (int max_E_cell;),
1027 (max_E_cell = (f.max_E_and_cell & 0x7FFFFFFFU) - 1;)
1028 );
1029
1030 CALORECGPU_CMC_LOAD
1031 (ClusterMaxCellEnergy,
1032 (ClusterMaxCellEnergyAndCell),
1033 (),
1034 (float max_E;),
1035 (max_E = __uint_as_float(f.max_E_and_cell >> 32U);)
1036 );
1037
1038 CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(ClusterSecondMaxCellEnergyAndCell, second_max_E_and_cell, secondMaxCellEnergyAndCell);
1039
1040 CALORECGPU_CMC_LOAD
1041 (ClusterCellWithSecondMaxEnergy,
1042 (ClusterSecondMaxCellEnergyAndCell),
1043 (),
1044 (int second_max_E_cell;),
1045 (second_max_E_cell = (f.second_max_E_and_cell & 0x7FFFFFFFU) - 1;)
1046 );
1047
1048 CALORECGPU_CMC_LOAD
1049 (ClusterSecondMaxCellEnergy,
1050 (ClusterSecondMaxCellEnergyAndCell),
1051 (),
1052 (float second_max_E;),
1053 (second_max_E = __uint_as_float(f.second_max_E_and_cell >> 32U);)
1054 );
1055
1056 CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(SumAbsEnergyNonMomentsBase, abs_energy_non_moments_base, sumAbsEnergyNonMoments);
1057 CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(SumAbsEnergyNonMomentsCorr, abs_energy_non_moments_corr, sumAbsEnergyNonMomentsAux);
1058
1059 CALORECGPU_CMC_LOAD
1060 (SumAbsEnergyNonMoments,
1061 (SumAbsEnergyNonMomentsBase, SumAbsEnergyNonMomentsCorr),
1062 (),
1063 (float abs_energy_non_moments;),
1064 (abs_energy_non_moments = f.abs_energy_non_moments_base + f.abs_energy_non_moments_corr;)
1065 );
1066
1067 CALORECGPU_CMC_LOAD
1068 (ReverseSumAbsEnergyNonMoments,
1069 (SumAbsEnergyNonMoments),
1070 (),
1071 (float rev_abs_energy_non_moments;),
1072 (rev_abs_energy_non_moments = 1.f / (f.abs_energy_non_moments != 0.f ? f.abs_energy_non_moments : 1.f);)
1073 );
1074
1075 CALORECGPU_CMC_LOAD_SIMPLE_MOMENT_INFO(CenterX, center_x, centerX);
1076 CALORECGPU_CMC_LOAD_SIMPLE_MOMENT_INFO(CenterY, center_y, centerY);
1077 CALORECGPU_CMC_LOAD_SIMPLE_MOMENT_INFO(CenterZ, center_z, centerZ);
1078
1079 CALORECGPU_CMC_LOAD_SIMPLE_CLUSTER_INFO(ClusterEnergy, cluster_energy, clusterEnergy);
1080
1081 CALORECGPU_CMC_LOAD
1082 (ReverseClusterEnergy,
1083 (ClusterEnergy),
1084 (),
1085 (float rev_cluster_energy;),
1086 (rev_cluster_energy = (f.cluster_energy != 0.f ? 1.f / f.cluster_energy : 1.f);)
1087 );
1088
1089 CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(MaxSignificanceAndSampling, max_sig_and_samp, maxSignificanceAndSampling);
1090
1091 CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(TimeNormalizationBase, time_norm_base, timeNormalization);
1092 CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(TimeNormalizationCorr, time_norm_corr, timeNormalizationAux);
1093
1094 CALORECGPU_CMC_LOAD
1095 (TimeNormalization,
1096 (TimeNormalizationBase, TimeNormalizationCorr),
1097 (),
1098 (float time_norm;),
1099 (time_norm = f.time_norm_base + f.time_norm_corr;)
1100 );
1101
1102 CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(ShowerAxisX, axis_x, showerAxisX);
1103 CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(ShowerAxisY, axis_y, showerAxisY);
1104 CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(ShowerAxisZ, axis_z, showerAxisZ);
1105
1106
1107 CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(ClusterMaxAndSecondMaxCellTogether, stored_max_and_second_max, maxAndSecondMaxCells);
1108
1109
1110 CALORECGPU_CMC_LOAD
1111 (ClusterMaxAndSecondMaxCell,
1112 (ClusterMaxAndSecondMaxCellTogether),
1113 (),
1114 (int max_cell, second_max_cell;),
1115 ( max_cell = f.stored_max_and_second_max >> 32U;
1116 second_max_cell = f.stored_max_and_second_max;
1117 )
1118 );
1119
1120 CALORECGPU_CMC_LOAD_SIMPLE_PER_SAMPLING_TEMPORARY_INFO(AbsoluteEnergyPerSample, sampling_normalization, absoluteEnergyPerSample);
1121
1122 CALORECGPU_CMC_LOAD
1123 (ReverseAbsoluteEnergyPerSample,
1124 (AbsoluteEnergyPerSample),
1125 (),
1126 (float rev_sampling_normalization;),
1127 (rev_sampling_normalization = 1.0f / (f.sampling_normalization != 0.f ? f.sampling_normalization : 1.0f);)
1128 );
1129
1130
1131 CALORECGPU_CMC_LOAD_SIMPLE_PER_SAMPLING_TEMPORARY_INFO(MaxEnergyAndCellPerSample, max_energy_and_cell, maxEnergyAndCellPerSample);
1132
1133 //+---------------------------------------------------+
1134 //| Others: |
1135 //+---------------------------------------------------+
1136
1137 //This one reads directly from the geometry array,
1138 //the calculation of the moments themselves probably benefits
1139 //from using the seedCellPhi temporary instead!
1140 CALORECGPU_CMC_LOAD
1141 (SeedCellGeometryPhi,
1142 (SeedCellIndex),
1143 (),
1144 (float seed_cell_phi_coordinate;),
1145 ( const int seed_cell_hash_ID = ( f.seed_cell >= 0 && f.seed_cell < CaloRecGPU::NCaloCells ?
1146 p.cell_info_arr->get_hash_ID(f.seed_cell, p.assume_complete_cells) : -1);
1147 seed_cell_phi_coordinate = (seed_cell_hash_ID >= 0 ? p.geometry->phi[seed_cell_hash_ID] : -999);
1148 )
1149 );
1150 }
1151
1152 //-------------------------------
1153 namespace ToCalculate
1154 {
1155
1156 struct Template
1157 {
1158 //A list of things to load (from memory)
1159 //so that we can load only once if we combine
1160 //several moments.
1161 using BeforeLoading = TypeList<>;
1162
1163 //T is a type that has at least
1164 //the variables specified in BeforeLoading
1165 template <class T>
1166 __device__ static void before(Parameters p,
1167 const T & data,
1168 const int cluster)
1169 {
1170 }
1171
1172 //CellLoading has access just to the cell index
1173 using CellLoading = TypeList<>;
1174 //ClusterLoading will have in the object passed to it
1175 //the cluster weight too.
1176 using ClusterLoading = TypeList<>;
1177
1178 //T is a type that has at least
1179 //the variables specified in CellLoading and ClusterLoading,
1180 //as well as the cluster weight.
1181 template <class T>
1182 __device__ static void per_cell(Parameters p,
1183 const T & data,
1184 const int cell,
1185 const int cluster)
1186 {
1187 }
1188
1189 using AfterLoading = TypeList<>;
1190
1191 //T is a type that has at least
1192 //the variables specified in AfterLoading
1193 template <class T>
1194 __device__ static void after(Parameters p,
1195 const T & data,
1196 const int cluster)
1197 {
1198 }
1199 };
1200
1201#define CALORECGPU_CMC_MOMENT_CALC(NAME, BEFORELOAD, BEFOREEXEC, CELLLOAD, CLUSTERLOAD, CELLEXEC, AFTERLOAD, AFTEREXEC) \
1202 struct NAME \
1203 { \
1204 using BeforeLoading = TypeList<CALORECGPU_CMC_EXPAND BEFORELOAD>; \
1205 template <class T> __device__ static void before(Parameters p, \
1206 const T & data, \
1207 const int cluster) \
1208 { CALORECGPU_CMC_EXPAND BEFOREEXEC } \
1209 using CellLoading = TypeList<CALORECGPU_CMC_EXPAND CELLLOAD>; \
1210 using ClusterLoading = TypeList<CALORECGPU_CMC_EXPAND CLUSTERLOAD>; \
1211 template <class T> __device__ static void per_cell(Parameters p, \
1212 const T & data, \
1213 const int cell, \
1214 const int cluster) \
1215 { CALORECGPU_CMC_EXPAND CELLEXEC } \
1216 using AfterLoading = TypeList<CALORECGPU_CMC_EXPAND AFTERLOAD>; \
1217 template <class T> __device__ static void after(Parameters p, \
1218 const T & data, \
1219 const int cluster) \
1220 { CALORECGPU_CMC_EXPAND AFTEREXEC } \
1221 }
1222
1223 //-------------------------------
1224
1225 //ISOLATION & ENG_FRAC_CORE SPECIAL CASING TO BE CONSIDERED LATER.
1226 //ALSO CENTER LAMBDA!
1227
1228 //+---------------------------------------------------+
1229 //| Regular cluster properties: |
1230 //+---------------------------------------------------+
1231
1232 CALORECGPU_CMC_MOMENT_CALC
1233 (ClusterEnergyEtaAndEt,
1234 (),
1235 (p.clusters_arr->clusterEnergy[cluster] = 0.f;
1236 CMCTemporaries::clusterEnergyAux(p.clusters_arr, cluster) = 0.f;
1237 p.clusters_arr->clusterEta[cluster] = 0.f;
1238 CMCTemporaries::clusterEtaAux(p.clusters_arr, cluster) = 0.f;),
1239 (ToLoad::CellEnergy, ToLoad::CellAbsEnergy, ToLoad::CellEta),
1240 (),
1241 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->clusterEnergy[cluster]),
1242 CMCTemporaries::clusterEnergyAux_ptr(p.clusters_arr, cluster),
1243 data.energy * data.weight);
1244 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->clusterEta[cluster]),
1245 CMCTemporaries::clusterEtaAux_ptr(p.clusters_arr, cluster),
1246 data.abs_energy * data.weight * data.eta);
1247 ),
1248 (ToLoad::SumAbsEnergyNonMoments, ToLoad::ReverseSumAbsEnergyNonMoments),
1249 (const float temp_E_main = p.clusters_arr->clusterEnergy[cluster];
1250 const float temp_E_corr = CMCTemporaries::clusterEnergyAux(p.clusters_arr, cluster);
1251
1252 const float temp_eta_main = p.clusters_arr->clusterEta[cluster];
1253 const float temp_eta_corr = CMCTemporaries::clusterEtaAux(p.clusters_arr, cluster);
1254
1255 const float temp_exp_1 = expf(temp_eta_main * data.rev_abs_energy_non_moments);
1256 const float temp_exp_2 = expf(temp_eta_corr * data.rev_abs_energy_non_moments);
1257
1258 const float temp_exp_mult = temp_exp_1 * temp_exp_2;
1259
1260 const float temp_numerator = 2.f * CaloRecGPU::Helpers::product_sum_cornea_harrison_tang(temp_E_main, temp_exp_mult,
1261 temp_E_corr, temp_exp_mult);
1262 const float temp_inv_denominator = 1.f/fmaf(temp_exp_mult, temp_exp_mult, 1.f);
1263
1264 const float temp_ET = temp_numerator * temp_inv_denominator;
1265
1266 const float temp_energy = temp_E_main + temp_E_corr;
1267
1268 p.clusters_arr->clusterEnergy[cluster] = temp_energy;
1269
1270 const float temp_eta = (temp_eta_main + temp_eta_corr) * data.rev_abs_energy_non_moments;
1271
1272 p.clusters_arr->clusterEta[cluster] = temp_eta * (data.abs_energy_non_moments != 0.f);
1273
1274 //const float temp_ET = temp_energy / coshf(abs(temp_eta));
1275
1276 p.clusters_arr->clusterEt[cluster] = temp_ET * (data.abs_energy_non_moments != 0.f);
1277 )
1278 );
1279
1280 CALORECGPU_CMC_MOMENT_CALC
1281 (ClusterPhi,
1282 (),
1283 (p.clusters_arr->clusterPhi[cluster] = 0.f;
1284 CMCTemporaries::clusterPhiAux(p.clusters_arr, cluster) = 0.f;
1285 ),
1286 (ToLoad::CellAbsEnergy, ToLoad::CellPhi),
1287 (ToLoad::SeedCellPhi),
1288 (const float phi_real = CaloRecGPU::Helpers::regularize_angle(data.phi, data.phi_0);
1289 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->clusterPhi[cluster]),
1290 CMCTemporaries::clusterPhiAux_ptr(p.clusters_arr, cluster),
1291 phi_real * data.abs_energy * data.weight);
1292 ),
1293 (ToLoad::SumAbsEnergyNonMoments, ToLoad::ReverseSumAbsEnergyNonMoments),
1294 (const float old_phi = p.clusters_arr->clusterPhi[cluster] + CMCTemporaries::clusterPhiAux(p.clusters_arr, cluster);
1295 p.clusters_arr->clusterPhi[cluster] = CaloRecGPU::Helpers::regularize_angle(old_phi *
1296 data.rev_abs_energy_non_moments, 0.f) *
1297 (data.abs_energy_non_moments != 0.f);
1298 )
1299 );
1300
1301 //+---------------------------------------------------+
1302 //| Cluster moments proper: |
1303 //+---------------------------------------------------+
1304
1305 CALORECGPU_CMC_MOMENT_CALC
1306 (AvgLArQ,
1307 (),
1308 (p.clusters_arr->moments.avgLArQ[cluster] = 0.f;
1309 CMCTemporaries::avgLArQAux(p.clusters_arr, cluster) = 0.f;
1310 ),
1311 (ToLoad::CellQualityProvenance, ToLoad::CellLArQCheck),
1312 (ToLoad::SquareWeightedEnergyOrNegative),
1313 (if (data.LArQ_cell_check)
1314 {
1315 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.avgLArQ[cluster]),
1316 CMCTemporaries::avgLArQAux_ptr(p.clusters_arr, cluster),
1317 data.square_w_E_or_neg * data.qp.quality());
1318 }
1319 ),
1320 (),
1321 (const float norm_LAr = CMCTemporaries::averageLArQNorm(p.clusters_arr, cluster) + CMCTemporaries::averageLArQNormAux(p.clusters_arr, cluster);
1322 const float rev_norm_LAr = 1.0f / (norm_LAr > 0.f ? norm_LAr : 1.0f);
1323 const float new_LArQ = p.clusters_arr->moments.avgLArQ[cluster] + CMCTemporaries::avgLArQAux(p.clusters_arr, cluster);
1324 p.clusters_arr->moments.avgLArQ[cluster] = new_LArQ * rev_norm_LAr;
1325 )
1326 );
1327
1328 CALORECGPU_CMC_MOMENT_CALC
1329 (AvgTileQ,
1330 (),
1331 (p.clusters_arr->moments.avgTileQ[cluster] = 0.f;
1332 CMCTemporaries::avgTileQAux(p.clusters_arr, cluster) = 0.f;
1333 ),
1334 (ToLoad::CellQualityProvenance, ToLoad::CellTileQCheck),
1335 (ToLoad::SquareWeightedEnergyOrNegative),
1336 (if (data.TileQ_cell_check)
1337 {
1338 const float max_quality = max((unsigned int) data.qp.tile_qual1(), (unsigned int) data.qp.tile_qual2());
1339
1340 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.avgTileQ[cluster]),
1341 CMCTemporaries::avgTileQAux_ptr(p.clusters_arr, cluster),
1342 data.square_w_E_or_neg * max_quality);
1343 }
1344 ),
1345 (),
1346 (const float norm_Tile = CMCTemporaries::averageTileQNorm(p.clusters_arr, cluster) + CMCTemporaries::averageTileQNormAux(p.clusters_arr, cluster);
1347 const float rev_norm_Tile = 1.0f / (norm_Tile > 0.f ? norm_Tile : 1.0f);
1348 const float new_TileQ = p.clusters_arr->moments.avgTileQ[cluster] + CMCTemporaries::avgTileQAux(p.clusters_arr, cluster);
1349 p.clusters_arr->moments.avgTileQ[cluster] = new_TileQ * rev_norm_Tile;
1350 )
1351 );
1352
1353 CALORECGPU_CMC_MOMENT_CALC
1354 (BadCellsCorrE,
1355 (),
1356 (p.clusters_arr->moments.badCellsCorrE[cluster] = 0.f;
1357 CMCTemporaries::badCellsCorrEAux(p.clusters_arr, cluster) = 0.f;
1358 ),
1359 (ToLoad::CellIsBad, ToLoad::CellEnergy),
1360 (ToLoad::WeightedEnergyOrNegative),
1361 (if (data.is_bad && data.energy != 0.f)
1362 {
1363 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.badCellsCorrE[cluster]),
1364 CMCTemporaries::badCellsCorrEAux_ptr(p.clusters_arr, cluster),
1365 data.weighted_energy_or_negative);
1366 }
1367 ),
1368 (),
1369 (p.clusters_arr->moments.badCellsCorrE[cluster] += CMCTemporaries::badCellsCorrEAux(p.clusters_arr, cluster);
1370 )
1371 );
1372
1373 CALORECGPU_CMC_MOMENT_CALC
1374 (BadLArQFrac,
1375 (),
1376 (p.clusters_arr->moments.badLArQFrac[cluster] = 0.f;
1377 CMCTemporaries::badLArQFracAux(p.clusters_arr, cluster) = 0.f;
1378 ),
1379 (ToLoad::CellQualityProvenance, ToLoad::CellLArQCheck),
1380 (ToLoad::WeightedEnergyOrNegative),
1381 (if (data.LArQ_cell_check && data.qp.quality() > p.opts->min_LAr_quality)
1382 {
1383 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.badLArQFrac[cluster]),
1384 CMCTemporaries::badLArQFracAux_ptr(p.clusters_arr, cluster),
1385 data.weighted_energy_or_negative);
1386 }
1387 ),
1388 (ToLoad::ReverseClusterEnergy),
1389 (const float new_badLArQFrac = p.clusters_arr->moments.badLArQFrac[cluster] + CMCTemporaries::badLArQFracAux(p.clusters_arr, cluster);
1390 p.clusters_arr->moments.badLArQFrac[cluster] = new_badLArQFrac * data.rev_cluster_energy;
1391 )
1392 );
1393
1394 CALORECGPU_CMC_MOMENT_CALC
1395 (CellSignificance,
1396 (),
1397 (),
1398 (),
1399 (),
1400 (),
1401 (ToLoad::MaxSignificanceAndSampling),
1402 (const float max_sig = __uint_as_float(data.max_sig_and_samp >> 32);
1403 p.clusters_arr->moments.cellSignificance[cluster] = max_sig * (data.max_sig_and_samp & 1 ? 1.f : -1.f);
1404 )
1405 );
1406
1407 CALORECGPU_CMC_MOMENT_CALC
1408 (CellSigSampling,
1409 (),
1410 (),
1411 (),
1412 (),
1413 (),
1414 (ToLoad::MaxSignificanceAndSampling),
1415 (const int max_samp = (data.max_sig_and_samp & 0xFFFFFFFEU) >> 1;
1416 p.clusters_arr->moments.cellSigSampling[cluster] = max_samp;
1417 )
1418 );
1419
1420
1421 CALORECGPU_CMC_MOMENT_CALC
1422 (CenterX,
1423 (),
1424 (p.clusters_arr->moments.centerX[cluster] = 0.f;
1425 CMCTemporaries::centerXAux(p.clusters_arr, cluster) = 0.f;
1426 ),
1427 (ToLoad::CellX),
1428 (ToLoad::WeightedEnergy),
1429 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.centerX[cluster]),
1430 CMCTemporaries::centerXAux_ptr(p.clusters_arr, cluster),
1431 data.x * data.weighted_energy);
1432 ),
1433 (ToLoad::ReverseSumEnergies),
1434 (const float new_value = p.clusters_arr->moments.centerX[cluster] + CMCTemporaries::centerXAux(p.clusters_arr, cluster);
1435 p.clusters_arr->moments.centerX[cluster] = new_value * data.rev_sum_energies;
1436 )
1437 );
1438 CALORECGPU_CMC_MOMENT_CALC
1439 (CenterY,
1440 (),
1441 (p.clusters_arr->moments.centerY[cluster] = 0.f;
1442 CMCTemporaries::centerYAux(p.clusters_arr, cluster) = 0.f;
1443 ),
1444 (ToLoad::CellY),
1445 (ToLoad::WeightedEnergy),
1446 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.centerY[cluster]),
1447 CMCTemporaries::centerYAux_ptr(p.clusters_arr, cluster),
1448 data.y * data.weighted_energy);
1449 ),
1450 (ToLoad::ReverseSumEnergies),
1451 (const float new_value = p.clusters_arr->moments.centerY[cluster] + CMCTemporaries::centerYAux(p.clusters_arr, cluster);
1452 p.clusters_arr->moments.centerY[cluster] = new_value * data.rev_sum_energies;
1453 )
1454 );
1455 CALORECGPU_CMC_MOMENT_CALC
1456 (CenterZ,
1457 (),
1458 (p.clusters_arr->moments.centerZ[cluster] = 0.f;
1459 CMCTemporaries::centerZAux(p.clusters_arr, cluster) = 0.f;
1460 ),
1461 (ToLoad::CellZ),
1462 (ToLoad::WeightedEnergy),
1463 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.centerZ[cluster]),
1464 CMCTemporaries::centerZAux_ptr(p.clusters_arr, cluster),
1465 data.z * data.weighted_energy);
1466 ),
1467 (ToLoad::ReverseSumEnergies),
1468 (const float new_value = p.clusters_arr->moments.centerZ[cluster] + CMCTemporaries::centerZAux(p.clusters_arr, cluster);
1469 p.clusters_arr->moments.centerZ[cluster] = new_value * data.rev_sum_energies;
1470 )
1471 );
1472
1473 CALORECGPU_CMC_MOMENT_CALC
1474 (EngFracEM,
1475 (),
1476 (p.clusters_arr->moments.engFracEM[cluster] = 0.f;
1477 CMCTemporaries::engFracEMAux(p.clusters_arr, cluster) = 0.f;),
1478 (ToLoad::CellSampling),
1479 (ToLoad::WeightedEnergy),
1480 ( if ( data.sampling == CaloSampling::EMB1 ||
1481 data.sampling == CaloSampling::EMB2 ||
1482 data.sampling == CaloSampling::EMB3 ||
1483 data.sampling == CaloSampling::EME1 ||
1484 data.sampling == CaloSampling::EME2 ||
1485 data.sampling == CaloSampling::EME3 ||
1486 data.sampling == CaloSampling::FCAL0 )
1487 {
1488 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.engFracEM[cluster]),
1489 CMCTemporaries::engFracEMAux_ptr(p.clusters_arr, cluster),
1490 data.weighted_energy);
1491 }
1492 ),
1493 (ToLoad::ReverseSumEnergies),
1494 (const float new_engFracEM = p.clusters_arr->moments.engFracEM[cluster] + CMCTemporaries::engFracEMAux(p.clusters_arr, cluster);
1495 p.clusters_arr->moments.engFracEM[cluster] = new_engFracEM * data.rev_sum_energies;
1496 )
1497 );
1498
1499 CALORECGPU_CMC_MOMENT_CALC
1500 (EngBadCells,
1501 (),
1502 (p.clusters_arr->moments.engBadCells[cluster] = 0.f;
1503 CMCTemporaries::engBadCellsAux(p.clusters_arr, cluster) = 0.f;),
1504 (ToLoad::CellIsBad),
1505 (ToLoad::WeightedEnergyOrNegative),
1506 (if (data.is_bad)
1507 {
1508 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.engBadCells[cluster]),
1509 CMCTemporaries::engBadCellsAux_ptr(p.clusters_arr, cluster),
1510 data.weighted_energy_or_negative);
1511 }
1512 ),
1513 (),
1514 (p.clusters_arr->moments.engBadCells[cluster] += CMCTemporaries::engBadCellsAux(p.clusters_arr, cluster);)
1515 );
1516
1517 CALORECGPU_CMC_MOMENT_CALC
1518 (EngFracMax,
1519 (),
1520 (),
1521 (),
1522 (),
1523 (),
1524 (ToLoad::ReverseSumEnergies, ToLoad::ClusterMaxCellEnergy),
1525 (p.clusters_arr->moments.engFracMax[cluster] = data.max_E * data.rev_sum_energies;)
1526 );
1527
1528 CALORECGPU_CMC_MOMENT_CALC
1529 (EngPosAndEngFracCore,
1530 (),
1531 (const float sum_energies = p.clusters_arr->moments.engPos[cluster] + CMCTemporaries::engPosAux(p.clusters_arr, cluster);
1532 p.clusters_arr->moments.engPos[cluster] = sum_energies;
1533 p.clusters_arr->moments.engFracCore[cluster] *= (sum_energies != 0.f ? 1.0f / sum_energies : 0.f);
1534 ),
1535 (),
1536 (),
1537 (),
1538 (),
1539 ()
1540 );
1541
1542 CALORECGPU_CMC_MOMENT_CALC
1543 (FirstEngDens,
1544 (),
1545 (p.clusters_arr->moments.firstEngDens[cluster] = 0.f;
1546 CMCTemporaries::firstEngDensAux(p.clusters_arr, cluster) = 0.f;),
1547 (ToLoad::CellVolume),
1548 (ToLoad::WeightedEnergy, ToLoad::WeightedEnergyOverVolume),
1549 (if (data.volume > 0.f)
1550 {
1551 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.firstEngDens[cluster]),
1552 CMCTemporaries::firstEngDensAux_ptr(p.clusters_arr, cluster),
1553 data.weighted_energy * data.w_E_over_V);
1554 }
1555 ),
1556 (ToLoad::ReverseEnergyDensityNormalization),
1557 (const float new_firstEngDens = p.clusters_arr->moments.firstEngDens[cluster] + CMCTemporaries::firstEngDensAux(p.clusters_arr, cluster);
1558 p.clusters_arr->moments.firstEngDens[cluster] = new_firstEngDens * data.rev_energy_density_norm;
1559 )
1560 );
1561
1562 CALORECGPU_CMC_MOMENT_CALC
1563 (FirstEta,
1564 (),
1565 (p.clusters_arr->moments.firstEta[cluster] = 0.f;
1566 CMCTemporaries::firstEtaAux(p.clusters_arr, cluster) = 0.f;
1567 ),
1568 (ToLoad::CellEta),
1569 (ToLoad::WeightedEnergy),
1570 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.firstEta[cluster]),
1571 CMCTemporaries::firstEtaAux_ptr(p.clusters_arr, cluster),
1572 data.weighted_energy * data.eta);
1573 ),
1574 (ToLoad::ReverseSumEnergies),
1575 (const float new_firstEta = p.clusters_arr->moments.firstEta[cluster] + CMCTemporaries::firstEtaAux(p.clusters_arr, cluster);
1576 p.clusters_arr->moments.firstEta[cluster] = new_firstEta * data.rev_sum_energies;
1577 )
1578 );
1579
1580 CALORECGPU_CMC_MOMENT_CALC
1581 (FirstPhi,
1582 (),
1583 (p.clusters_arr->moments.firstPhi[cluster] = 0.f;
1584 CMCTemporaries::firstPhiAux(p.clusters_arr, cluster) = 0.f;
1585 ),
1586 (ToLoad::CellPhi),
1587 (ToLoad::SeedCellPhi, ToLoad::WeightedEnergy),
1588 (const float phi_real = CaloRecGPU::Helpers::regularize_angle(data.phi, data.phi_0);
1589 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.firstPhi[cluster]),
1590 CMCTemporaries::firstPhiAux_ptr(p.clusters_arr, cluster),
1591 data.weighted_energy * phi_real);
1592 ),
1593 (ToLoad::ReverseSumEnergies),
1594 (const float new_firstPhi = p.clusters_arr->moments.firstPhi[cluster] + CMCTemporaries::firstPhiAux(p.clusters_arr, cluster);
1595 p.clusters_arr->moments.firstPhi[cluster] = CaloRecGPU::Helpers::regularize_angle(new_firstPhi * data.rev_sum_energies);
1596 )
1597 );
1598
1599 CALORECGPU_CMC_MOMENT_CALC
1600 (Lateral,
1601 (),
1602 (p.clusters_arr->moments.lateral[cluster] = 0.f;
1603 CMCTemporaries::lateralAux(p.clusters_arr, cluster) = 0.f;
1604 ),
1605 (),
1606 (ToLoad::WeightedEnergy, ToLoad::ClusterMaxAndSecondMaxCell, ToLoad::R),
1607 (if (cell != data.max_cell && cell != data.second_max_cell)
1608 {
1609 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.lateral[cluster]),
1610 CMCTemporaries::lateralAux_ptr(p.clusters_arr, cluster),
1611 data.weighted_energy * data.r * data.r);
1612 }
1613 ),
1614 (),
1615 (const float new_lateral = p.clusters_arr->moments.lateral[cluster] + CMCTemporaries::lateralAux(p.clusters_arr, cluster);
1616 const float new_norm = CMCTemporaries::lateralNormalization(p.clusters_arr, cluster) +
1617 CMCTemporaries::lateralNormalizationAux(p.clusters_arr, cluster);
1618 p.clusters_arr->moments.lateral[cluster] = new_lateral / (new_norm != 0.f ? new_norm : 1.f);
1619 )
1620 );
1621
1622 CALORECGPU_CMC_MOMENT_CALC
1623 (Longitudinal,
1624 (),
1625 (p.clusters_arr->moments.longitudinal[cluster] = 0.f;
1626 CMCTemporaries::longitudinalAux(p.clusters_arr, cluster) = 0.f;
1627 ),
1628 (),
1629 (ToLoad::WeightedEnergy, ToLoad::ClusterMaxAndSecondMaxCell, ToLoad::Lambda),
1630 (if (cell != data.max_cell && cell != data.second_max_cell)
1631 {
1632 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.longitudinal[cluster]),
1633 CMCTemporaries::longitudinalAux_ptr(p.clusters_arr, cluster),
1634 data.weighted_energy * data.lambda * data.lambda);
1635 }
1636 ),
1637 (),
1638 (const float new_longitudinal = p.clusters_arr->moments.longitudinal[cluster] + CMCTemporaries::longitudinalAux(p.clusters_arr, cluster);
1639 const float new_norm = CMCTemporaries::longitudinalNormalization(p.clusters_arr, cluster) +
1640 CMCTemporaries::longitudinalNormalizationAux(p.clusters_arr, cluster);
1641 p.clusters_arr->moments.longitudinal[cluster] = new_longitudinal / (new_norm != 0.f ? new_norm : 1.f);
1642 )
1643 );
1644
1645 CALORECGPU_CMC_MOMENT_CALC
1646 (Mass,
1647 (),
1648 (),
1649 (),
1650 (),
1651 (),
1652 (ToLoad::SumEnergies),
1653 (const float mx = CMCTemporaries::mX(p.clusters_arr, cluster) + CMCTemporaries::mXAux(p.clusters_arr, cluster);
1654 const float my = CMCTemporaries::mY(p.clusters_arr, cluster) + CMCTemporaries::mYAux(p.clusters_arr, cluster);
1655 const float mz = CMCTemporaries::mZ(p.clusters_arr, cluster) + CMCTemporaries::mZAux(p.clusters_arr, cluster);
1656
1657 const float v_1 = mx * mx;
1658 const float v_2 = my * my;
1659 const float v_3 = mz * mz;
1660 const float v_4 = data.sum_energies * data.sum_energies;
1661
1662 const float c_1 = fmaf(mx, mx, -v_1);
1663 const float c_2 = fmaf(my, my, -v_2);
1664 const float c_3 = fmaf(mz, mz, -v_3);
1665 const float c_4 = fmaf(data.sum_energies, data.sum_energies, -v_4);
1666
1667 const float sq_mass = CaloRecGPU::Helpers::sum_kahan_babushka_neumaier(v_4, -v_1, -v_2, -v_3, c_4, -c_1, -c_2, -c_3);
1668
1669 p.clusters_arr->moments.mass[cluster] = sqrtf(fabsf(sq_mass)) * ((sq_mass > 0.f) - (sq_mass < 0.f));
1670 )
1671 );
1672
1673 CALORECGPU_CMC_MOMENT_CALC
1674 (NBadCells,
1675 (),
1676 (p.clusters_arr->moments.nBadCells[cluster] = 0;),
1677 (ToLoad::CellIsBad),
1678 (),
1679 (if (data.is_bad)
1680 {
1681 atomicAdd(&(p.clusters_arr->moments.nBadCells[cluster]), 1);
1682 }),
1683 (),
1684 ()
1685 );
1686
1687 CALORECGPU_CMC_MOMENT_CALC
1688 (NBadCellsCorr,
1689 (),
1690 (p.clusters_arr->moments.nBadCellsCorr[cluster] = 0;),
1691 (ToLoad::CellIsBad, ToLoad::CellEnergy),
1692 (),
1693 (if (data.is_bad && data.energy != 0.f)
1694 {
1695 atomicAdd(&(p.clusters_arr->moments.nBadCellsCorr[cluster]), 1);
1696 }
1697 ),
1698 (),
1699 ()
1700 );
1701
1702 CALORECGPU_CMC_MOMENT_CALC
1703 (NExtraCellSampling,
1704 (),
1705 (p.clusters_arr->moments.nExtraCellSampling[cluster] = 0;),
1706 (ToLoad::CellSampling, ToLoad::CellEta),
1707 (),
1708 (if (data.sampling == CaloSampling::EME2 && fabsf(data.eta) > p.opts->eta_inner_wheel)
1709 {
1710 atomicAdd(&(p.clusters_arr->moments.nExtraCellSampling[cluster]), 1);
1711 }
1712 ),
1713 (),
1714 ()
1715 );
1716
1717 CALORECGPU_CMC_MOMENT_CALC
1718 (PTD,
1719 (),
1720 (p.clusters_arr->moments.PTD[cluster] = 0.f;
1721 CMCTemporaries::PTDAux(p.clusters_arr, cluster) = 0.f;
1722 ),
1723 (),
1724 (ToLoad::SquareWeightedEnergy),
1725 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.PTD[cluster]),
1726 CMCTemporaries::PTDAux_ptr(p.clusters_arr, cluster),
1727 data.square_w_E);
1728
1729 //Comment on there:
1730 //
1731 // +--------------- begin comment on there ---------------+
1732 // | |
1733 // | do not convert to pT since clusters are small and |
1734 // | there is virtually no difference and cosh just costs |
1735 // | time ... |
1736 // | |
1737 // +---------------- end comment on there ----------------+
1738 //
1739 //So maybe we could change this here?
1740 ),
1741 (ToLoad::SumEnergies),
1742 (const float new_PTD = p.clusters_arr->moments.PTD[cluster] + CMCTemporaries::PTDAux(p.clusters_arr, cluster);
1743 p.clusters_arr->moments.PTD[cluster] = 1.0f / ((data.sum_energies > 0.f ? data.sum_energies : 1.f) * rsqrtf(new_PTD));
1744 //See before: maybe to be revised?
1745 )
1746 );
1747
1748 CALORECGPU_CMC_MOMENT_CALC
1749 (SecondEngDens,
1750 (),
1751 (p.clusters_arr->moments.secondEngDens[cluster] = 0.f;
1752 CMCTemporaries::secondEngDensAux(p.clusters_arr, cluster) = 0.f;
1753 ),
1754 (ToLoad::CellVolume),
1755 (ToLoad::WeightedEnergy, ToLoad::WeightedEnergyOverVolume),
1756 (if (data.volume > 0.f)
1757 {
1758 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.secondEngDens[cluster]),
1759 CMCTemporaries::secondEngDensAux_ptr(p.clusters_arr, cluster),
1760 data.weighted_energy * data.w_E_over_V * data.w_E_over_V);
1761 }
1762 ),
1763 (ToLoad::ReverseEnergyDensityNormalization),
1764 (const float new_secondEngDens = p.clusters_arr->moments.secondEngDens[cluster] + CMCTemporaries::secondEngDensAux(p.clusters_arr, cluster);
1765 p.clusters_arr->moments.secondEngDens[cluster] = new_secondEngDens * data.rev_energy_density_norm;)
1766 );
1767
1768 CALORECGPU_CMC_MOMENT_CALC
1769 (SecondLambda,
1770 (),
1771 (p.clusters_arr->moments.secondLambda[cluster] = 0.f;
1772 CMCTemporaries::secondLambdaAux(p.clusters_arr, cluster) = 0.f;
1773 ),
1774 (),
1775 (ToLoad::WeightedEnergy, ToLoad::Lambda),
1776 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.secondLambda[cluster]),
1777 CMCTemporaries::secondLambdaAux_ptr(p.clusters_arr, cluster),
1778 data.weighted_energy * data.lambda * data.lambda);
1779 ),
1780 (ToLoad::ReverseSumEnergies),
1781 (const float new_secondLambda = p.clusters_arr->moments.secondLambda[cluster] + CMCTemporaries::secondLambdaAux(p.clusters_arr, cluster);
1782 p.clusters_arr->moments.secondLambda[cluster] = new_secondLambda * data.rev_sum_energies;
1783 )
1784 );
1785
1786 CALORECGPU_CMC_MOMENT_CALC
1787 (SecondR,
1788 (),
1789 (p.clusters_arr->moments.secondR[cluster] = 0.f;
1790 CMCTemporaries::secondRAux(p.clusters_arr, cluster) = 0.f;
1791 ),
1792 (),
1793 (ToLoad::WeightedEnergy, ToLoad::R),
1794 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.secondR[cluster]),
1795 CMCTemporaries::secondRAux_ptr(p.clusters_arr, cluster),
1796 data.weighted_energy * data.r * data.r);
1797 ),
1798 (ToLoad::ReverseSumEnergies),
1799 (const float new_secondR = p.clusters_arr->moments.secondR[cluster] + CMCTemporaries::secondRAux(p.clusters_arr, cluster);
1800 p.clusters_arr->moments.secondR[cluster] = new_secondR * data.rev_sum_energies;
1801 )
1802 );
1803
1804
1805 CALORECGPU_CMC_MOMENT_CALC
1806 (Significance,
1807 (),
1808 (p.clusters_arr->moments.significance[cluster] = 0.f;
1809 CMCTemporaries::significanceAux(p.clusters_arr, cluster) = 0.f;
1810 ),
1811 (ToLoad::CellNoise),
1812 (),
1813 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.significance[cluster]),
1814 CMCTemporaries::significanceAux_ptr(p.clusters_arr, cluster),
1815 data.noise * data.noise);
1816 ),
1817 (ToLoad::ClusterEnergy),
1818 (const float prev_v = p.clusters_arr->moments.significance[cluster] + CMCTemporaries::significanceAux(p.clusters_arr, cluster);
1819 p.clusters_arr->moments.significance[cluster] = (prev_v > 0.f ? data.cluster_energy * rsqrtf(prev_v) : 0.f);)
1820 );
1821
1822 CALORECGPU_CMC_MOMENT_CALC
1823 (TimeAndSecondTime,
1824 (),
1825 (p.clusters_arr->moments.time[cluster] = 0.f;
1826 CMCTemporaries::timeAux(p.clusters_arr, cluster) = 0.f;
1827 p.clusters_arr->moments.secondTime[cluster] = 0.f;
1828 CMCTemporaries::secondTimeAux(p.clusters_arr, cluster) = 0.f;
1829 ),
1830 (ToLoad::CellTime, ToLoad::CellTimeMomentsCheck),
1831 (ToLoad::SquaredWeightedNonMomentsEnergy),
1832 (if (data.time_moments_check)
1833 {
1834 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.time[cluster]),
1835 CMCTemporaries::timeAux_ptr(p.clusters_arr, cluster),
1836 data.time * data.squared_normE);
1837 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.secondTime[cluster]),
1838 CMCTemporaries::secondTimeAux_ptr(p.clusters_arr, cluster),
1839 data.time * data.time * data.squared_normE);
1840 }),
1841 (ToLoad::TimeNormalization),
1842 (if (data.time_norm != 0.f)
1843 {
1844 const float real_norm = 1.0f / data.time_norm;
1845 const float time = (p.clusters_arr->moments.time[cluster] + CMCTemporaries::timeAux(p.clusters_arr, cluster))
1846 * real_norm;
1847 const float second_sum = p.clusters_arr->moments.secondTime[cluster] + CMCTemporaries::secondTimeAux(p.clusters_arr, cluster);
1848 p.clusters_arr->moments.time[cluster] = time;
1849 p.clusters_arr->moments.secondTime[cluster] = CaloRecGPU::Helpers::product_sum_cornea_harrison_tang(second_sum, real_norm, -time, time);
1850 }
1851 else
1852 {
1853 p.clusters_arr->moments.time[cluster] = 0.f;
1854 p.clusters_arr->moments.secondTime[cluster] = 0.f;
1855 }
1856 )
1857 );
1858
1859 //+---------------------------------------------------+
1860 //| Cluster per-sampling moments: |
1861 //+---------------------------------------------------+
1862
1863 template <int num, int delta = 0>
1864 CALORECGPU_CMC_MOMENT_CALC
1865 (EnergyPerSampleSeveral,
1866 (ToLoad::SamplingFromMomentIndex),
1867 (const int offset = data.sampling * num + delta;
1868 for (int i = 0; i < num; ++i)
1869 {
1870 p.clusters_arr->moments.energyPerSample[offset + i]
1871 [cluster] = 0.f;
1872 CMCTemporaries::energyPerSampleAux(p.clusters_arr, offset + i, cluster) = 0.f;
1873 }
1874 ),
1875 (ToLoad::CellEnergy, ToLoad::CellSampling),
1876 (),
1877 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.energyPerSample[data.sampling][cluster]),
1878 CMCTemporaries::energyPerSampleAux_ptr(p.clusters_arr, data.sampling, cluster),
1879 data.energy * data.weight);
1880 ),
1881 (ToLoad::SamplingFromMomentIndex),
1882 (const int offset = data.sampling * num + delta;
1883 for (int i = 0; i < num; ++i)
1884 {
1885 p.clusters_arr->moments.energyPerSample[offset + i]
1886 [cluster] += CMCTemporaries::energyPerSampleAux(p.clusters_arr, offset + i, cluster);
1887 }
1888 )
1889 );
1890
1891 using EnergyPerSample = EnergyPerSampleSeveral<1, 0>;
1892
1893
1894 template <int num, int delta = 0>
1895 CALORECGPU_CMC_MOMENT_CALC
1896 (EtaPerSampleSeveral,
1897 (ToLoad::SamplingFromMomentIndex),
1898 (const int offset = data.sampling * num + delta;
1899 for (int i = 0; i < num; ++i)
1900 {
1901 p.clusters_arr->moments.etaPerSample[offset + i]
1902 [cluster] = 0.f;
1903 CMCTemporaries::etaPerSampleAux(p.clusters_arr, offset + i, cluster) = 0.f;
1904 }
1905 ),
1906 (ToLoad::CellSampling, ToLoad::CellAbsEnergy, ToLoad::CellEta),
1907 (),
1908 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.etaPerSample[data.sampling][cluster]),
1909 CMCTemporaries::etaPerSampleAux_ptr(p.clusters_arr, data.sampling, cluster),
1910 data.abs_energy * data.weight * data.eta);
1911 ),
1912 (ToLoad::SamplingFromMomentIndex/*, ToLoad::ReverseAbsoluteEnergyPerSample*/),
1913 (const int offset = data.sampling * num + delta;
1914 for (int i = 0; i < num; ++i)
1915 {
1916 const int idx = offset + i;
1917 const float normalization = CMCTemporaries::absoluteEnergyPerSample(p.clusters_arr, idx, cluster) +
1918 CMCTemporaries::absoluteEnergyPerSampleAux(p.clusters_arr, idx, cluster);
1919 const float rev_normalization = 1.0f / (normalization != 0.f ? normalization : 1.0f);
1920 const float new_eta = p.clusters_arr->moments.etaPerSample[idx][cluster] + CMCTemporaries::etaPerSampleAux(p.clusters_arr, idx, cluster);
1921 p.clusters_arr->moments.etaPerSample[idx][cluster] = new_eta * rev_normalization;
1922 }
1923 )
1924 );
1925
1926 using EtaPerSample = EtaPerSampleSeveral<1, 0>;
1927
1928 template <int num, int delta = 0>
1929 CALORECGPU_CMC_MOMENT_CALC
1930 (NCellSamplingSeveral,
1931 (ToLoad::SamplingFromMomentIndex),
1932 (const int offset = data.sampling * num + delta;
1933 for (int i = 0; i < num; ++i)
1934 {
1935 p.clusters_arr->moments.nCellSampling[offset + i]
1936 [cluster] = 0;
1937 }
1938 ),
1939 (ToLoad::CellSampling),
1940 (),
1941 (atomicAdd(&(p.clusters_arr->moments.nCellSampling[data.sampling][cluster]), 1);),
1942 (),
1943 ()
1944 );
1945
1946 using NCellSampling = NCellSamplingSeveral<1, 0>;
1947
1948 template <int num, int delta = 0>
1949 CALORECGPU_CMC_MOMENT_CALC
1950 (PhiPerSampleSeveral,
1951 (ToLoad::SamplingFromMomentIndex),
1952 (const int offset = data.sampling * num + delta;
1953 for (int i = 0; i < num; ++i)
1954 {
1955 p.clusters_arr->moments.phiPerSample[offset + i]
1956 [cluster] = 0.f;
1957 CMCTemporaries::phiPerSampleAux(p.clusters_arr, offset + i, cluster) = 0.f;
1958 }
1959 ),
1960 (ToLoad::CellSampling, ToLoad::CellAbsEnergy, ToLoad::CellPhi),
1961 (ToLoad::SeedCellPhi),
1962 (const float phi_real = CaloRecGPU::Helpers::regularize_angle(data.phi, data.phi_0);
1963 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(&(p.clusters_arr->moments.phiPerSample[data.sampling][cluster]),
1964 CMCTemporaries::phiPerSampleAux_ptr(p.clusters_arr, data.sampling, cluster),
1965 data.abs_energy * data.weight * phi_real);
1966 ),
1967 (ToLoad::SamplingFromMomentIndex/*, ToLoad::ReverseAbsoluteEnergyPerSample*/),
1968 (const int offset = data.sampling * num + delta;
1969 for (int i = 0; i < num; ++i)
1970 {
1971 const int idx = offset + i;
1972 const float normalization = CMCTemporaries::absoluteEnergyPerSample(p.clusters_arr, idx, cluster) +
1973 CMCTemporaries::absoluteEnergyPerSampleAux(p.clusters_arr, idx, cluster);
1974 const float rev_normalization = 1.0f / (normalization != 0.f ? normalization : 1.0f);
1975 const float new_phi = p.clusters_arr->moments.phiPerSample[idx][cluster] + CMCTemporaries::phiPerSampleAux(p.clusters_arr, idx, cluster);
1976 p.clusters_arr->moments.phiPerSample[idx][cluster] = CaloRecGPU::Helpers::regularize_angle(new_phi * rev_normalization, 0.f);
1977 }
1978 )
1979 );
1980
1981 using PhiPerSample = PhiPerSampleSeveral<1, 0>;
1982
1983 //+---------------------------------------------------+
1984 //| Temporary Values: |
1985 //+---------------------------------------------------+
1986
1987
1988 CALORECGPU_CMC_MOMENT_CALC
1989 (AverageLArQNormalization,
1990 (),
1991 (CMCTemporaries::averageLArQNorm(p.clusters_arr, cluster) = 0.f;
1992 CMCTemporaries::averageLArQNormAux(p.clusters_arr, cluster) = 0.f;),
1993 (ToLoad::CellLArQCheck),
1994 (ToLoad::SquareWeightedEnergyOrNegative),
1995 (if (data.LArQ_cell_check)
1996 {
1997 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::averageLArQNorm_ptr(p.clusters_arr, cluster),
1998 CMCTemporaries::averageLArQNormAux_ptr(p.clusters_arr, cluster),
1999 data.square_w_E_or_neg);
2000 }
2001 ),
2002 (),
2003 ()
2004 );
2005
2006 CALORECGPU_CMC_MOMENT_CALC
2007 (AverageTileQNormalization,
2008 (),
2009 (CMCTemporaries::averageTileQNorm(p.clusters_arr, cluster) = 0.f;
2010 CMCTemporaries::averageTileQNormAux(p.clusters_arr, cluster) = 0.f;),
2011 (ToLoad::CellQualityProvenance, ToLoad::CellTileQCheck),
2012 (ToLoad::SquareWeightedEnergyOrNegative),
2013 (if (data.TileQ_cell_check)
2014 {
2015 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::averageTileQNorm_ptr(p.clusters_arr, cluster),
2016 CMCTemporaries::averageTileQNormAux_ptr(p.clusters_arr, cluster),
2017 data.square_w_E_or_neg);
2018 }
2019 ),
2020 (),
2021 ()
2022 );
2023
2024 CALORECGPU_CMC_MOMENT_CALC
2025 (EnergyDensityNormalization,
2026 (),
2027 (CMCTemporaries::energyDensityNormalization(p.clusters_arr, cluster) = 0.f;
2028 CMCTemporaries::energyDensityNormalizationAux(p.clusters_arr, cluster) = 0.f;
2029 ),
2030 (ToLoad::CellVolume),
2031 (ToLoad::WeightedEnergy),
2032 (if (data.volume > 0.f)
2033 {
2034 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::energyDensityNormalization_ptr(p.clusters_arr, cluster),
2035 CMCTemporaries::energyDensityNormalizationAux_ptr(p.clusters_arr, cluster),
2036 data.weighted_energy);
2037 }
2038 ),
2039 (),
2040 ()
2041 );
2042
2043 CALORECGPU_CMC_MOMENT_CALC
2044 (FirstAndSecondMaxEnergyAndCell,
2045 (),
2046 (CMCTemporaries::maxCellEnergyAndCell(p.clusters_arr, cluster) = 0ULL;
2047 CMCTemporaries::secondMaxCellEnergyAndCell(p.clusters_arr, cluster) = 0ULL;
2048 ),
2049 (),
2050 (ToLoad::WeightedEnergy),
2051 (if (data.weighted_energy > 0)
2052 {
2053 unsigned long long int energy_and_cell = __float_as_uint(data.weighted_energy);
2054 //Energy is positive, so no need to switch to total ordering...
2055 energy_and_cell = (energy_and_cell << 32) | (cell + 1);
2056 const unsigned long long int old_enc = atomicMax(&(CMCTemporaries::maxCellEnergyAndCell(p.clusters_arr, cluster)), energy_and_cell);
2057 atomicMax(&(CMCTemporaries::secondMaxCellEnergyAndCell(p.clusters_arr, cluster)), min(old_enc, energy_and_cell));
2058 }
2059 ),
2060 (),
2061 ()
2062 );
2063
2064
2065 CALORECGPU_CMC_MOMENT_CALC
2066 (LateralNormalization,
2067 (),
2068 (CMCTemporaries::lateralNormalization(p.clusters_arr, cluster) = 0.f;
2069 CMCTemporaries::lateralNormalizationAux(p.clusters_arr, cluster) = 0.f;
2070 ),
2071 (),
2072 (ToLoad::WeightedEnergy, ToLoad::ClusterMaxAndSecondMaxCell, ToLoad::R),
2073 (const float real_r = (cell != data.max_cell && cell != data.second_max_cell) ?
2074 data.r : max(data.r, p.opts->min_r_lateral);
2075
2076 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::lateralNormalization_ptr(p.clusters_arr, cluster),
2077 CMCTemporaries::lateralNormalizationAux_ptr(p.clusters_arr, cluster),
2078 data.weighted_energy * real_r * real_r);
2079 ),
2080 (),
2081 ()
2082 );
2083
2084 CALORECGPU_CMC_MOMENT_CALC
2085 (LongitudinalNormalization,
2086 (),
2087 (CMCTemporaries::longitudinalNormalization(p.clusters_arr, cluster) = 0.f;
2088 CMCTemporaries::longitudinalNormalizationAux(p.clusters_arr, cluster) = 0.f;
2089 ),
2090 (),
2091 (ToLoad::WeightedEnergy, ToLoad::ClusterMaxAndSecondMaxCell, ToLoad::Lambda),
2092 (const float real_lambda = (cell != data.max_cell && cell != data.second_max_cell) ?
2093 data.lambda : max(data.lambda, p.opts->min_l_longitudinal);
2094
2095 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::longitudinalNormalization_ptr(p.clusters_arr, cluster),
2096 CMCTemporaries::longitudinalNormalizationAux_ptr(p.clusters_arr, cluster),
2097 data.weighted_energy * real_lambda * real_lambda);
2098 ),
2099 (),
2100 ()
2101 );
2102
2103 CALORECGPU_CMC_MOMENT_CALC
2104 (Matrix00,
2105 (),
2106 (CMCTemporaries::matrix00(p.clusters_arr, cluster) = 0.f;
2107 CMCTemporaries::matrix00Aux(p.clusters_arr, cluster) = 0.f;
2108 ),
2109 (ToLoad::CellX),
2110 (ToLoad::CenterX, ToLoad::SquareWeightedEnergy),
2111 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::matrix00_ptr(p.clusters_arr, cluster),
2112 CMCTemporaries::matrix00Aux_ptr(p.clusters_arr, cluster),
2113 data.square_w_E * (data.x - data.center_x) * (data.x - data.center_x));
2114 ),
2115 (),
2116 ()
2117 );
2118 CALORECGPU_CMC_MOMENT_CALC
2119 (Matrix10,
2120 (),
2121 (CMCTemporaries::matrix10(p.clusters_arr, cluster) = 0.f;
2122 CMCTemporaries::matrix10Aux(p.clusters_arr, cluster) = 0.f;
2123 ),
2124 (ToLoad::CellX, ToLoad::CellY),
2125 (ToLoad::CenterX, ToLoad::CenterY, ToLoad::SquareWeightedEnergy),
2126 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::matrix10_ptr(p.clusters_arr, cluster),
2127 CMCTemporaries::matrix10Aux_ptr(p.clusters_arr, cluster),
2128 data.square_w_E * (data.x - data.center_x) * (data.y - data.center_y));
2129 ),
2130 (),
2131 ()
2132 );
2133 CALORECGPU_CMC_MOMENT_CALC
2134 (Matrix20,
2135 (),
2136 (CMCTemporaries::matrix20(p.clusters_arr, cluster) = 0.f;
2137 CMCTemporaries::matrix20Aux(p.clusters_arr, cluster) = 0.f;
2138 ),
2139 (ToLoad::CellX, ToLoad::CellZ),
2140 (ToLoad::CenterX, ToLoad::CenterZ, ToLoad::SquareWeightedEnergy),
2141 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::matrix20_ptr(p.clusters_arr, cluster),
2142 CMCTemporaries::matrix20Aux_ptr(p.clusters_arr, cluster),
2143 data.square_w_E * (data.x - data.center_x) * (data.z - data.center_z));
2144 ),
2145 (),
2146 ()
2147 );
2148 CALORECGPU_CMC_MOMENT_CALC
2149 (Matrix11,
2150 (),
2151 (CMCTemporaries::matrix11(p.clusters_arr, cluster) = 0.f;
2152 CMCTemporaries::matrix11Aux(p.clusters_arr, cluster) = 0.f;
2153 ),
2154 (ToLoad::CellY),
2155 (ToLoad::CenterY, ToLoad::SquareWeightedEnergy),
2156 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::matrix11_ptr(p.clusters_arr, cluster),
2157 CMCTemporaries::matrix11Aux_ptr(p.clusters_arr, cluster),
2158 data.square_w_E * (data.y - data.center_y) * (data.y - data.center_y));
2159 ),
2160 (),
2161 ()
2162 );
2163 CALORECGPU_CMC_MOMENT_CALC
2164 (Matrix21,
2165 (),
2166 (CMCTemporaries::matrix21(p.clusters_arr, cluster) = 0.f;
2167 CMCTemporaries::matrix21Aux(p.clusters_arr, cluster) = 0.f;
2168 ),
2169 (ToLoad::CellY, ToLoad::CellZ),
2170 (ToLoad::CenterY, ToLoad::CenterZ, ToLoad::SquareWeightedEnergy),
2171 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::matrix21_ptr(p.clusters_arr, cluster),
2172 CMCTemporaries::matrix21Aux_ptr(p.clusters_arr, cluster),
2173 data.square_w_E * (data.y - data.center_y) * (data.z - data.center_z));
2174 ),
2175 (),
2176 ()
2177 );
2178 CALORECGPU_CMC_MOMENT_CALC
2179 (Matrix22,
2180 (),
2181 (CMCTemporaries::matrix22(p.clusters_arr, cluster) = 0.f;
2182 CMCTemporaries::matrix22Aux(p.clusters_arr, cluster) = 0.f;
2183 ),
2184 (ToLoad::CellZ),
2185 (ToLoad::CenterZ, ToLoad::SquareWeightedEnergy),
2186 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::matrix22_ptr(p.clusters_arr, cluster),
2187 CMCTemporaries::matrix22Aux_ptr(p.clusters_arr, cluster),
2188 data.square_w_E * (data.z - data.center_z) * (data.z - data.center_z));
2189 ),
2190 (),
2191 ()
2192 );
2193
2194 CALORECGPU_CMC_MOMENT_CALC
2195 (MaxAndSecondMaxCells,
2196 (),
2197 (),
2198 (),
2199 (),
2200 (),
2201 (ToLoad::ClusterCellWithMaxEnergy, ToLoad::ClusterCellWithSecondMaxEnergy),
2202 (unsigned long long to_store = data.max_E_cell;
2203 to_store = (to_store << 32u) | static_cast<unsigned int>(data.second_max_E_cell);
2204 CMCTemporaries::maxAndSecondMaxCells(p.clusters_arr, cluster) = to_store;
2205 )
2206 );
2207
2208 CALORECGPU_CMC_MOMENT_CALC
2209 (MaxSignificanceAndSampling,
2210 (),
2211 (CMCTemporaries::maxSignificanceAndSampling(p.clusters_arr, cluster) = 0ULL;),
2212 (ToLoad::CellSampling, ToLoad::CellNoise),
2213 (ToLoad::WeightedEnergyOrNegative),
2214 (const float max_sig = data.noise > 0.f ? data.weighted_energy_or_negative / data.noise : 0.f;
2215 unsigned long long int max_S_and_S = __float_as_uint(fabsf(max_sig));
2216 max_S_and_S = (max_S_and_S << 32) | (static_cast<unsigned long long int>(data.sampling) << 1) | (max_sig > 0.f);
2217 atomicMax(&(CMCTemporaries::maxSignificanceAndSampling(p.clusters_arr, cluster)), max_S_and_S);),
2218 (),
2219 ()
2220 );
2221
2222 CALORECGPU_CMC_MOMENT_CALC
2223 (MX,
2224 (),
2225 (CMCTemporaries::mX(p.clusters_arr, cluster) = 0.f;
2226 CMCTemporaries::mXAux(p.clusters_arr, cluster) = 0.f;),
2227 (ToLoad::CellX),
2228 (ToLoad::WeightedCellPositionNormalization),
2229 (const float mx = data.w_E_r_dir * data.x;
2230 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::mX_ptr(p.clusters_arr, cluster),
2231 CMCTemporaries::mXAux_ptr(p.clusters_arr, cluster),
2232 mx);
2233 ),
2234 (),
2235 ()
2236 );
2237 CALORECGPU_CMC_MOMENT_CALC
2238 (MY,
2239 (),
2240 (CMCTemporaries::mY(p.clusters_arr, cluster) = 0.f;
2241 CMCTemporaries::mYAux(p.clusters_arr, cluster) = 0.f;),
2242 (ToLoad::CellY),
2243 (ToLoad::WeightedCellPositionNormalization),
2244 (const float my = data.w_E_r_dir * data.y;
2245 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::mY_ptr(p.clusters_arr, cluster),
2246 CMCTemporaries::mYAux_ptr(p.clusters_arr, cluster),
2247 my);
2248 ),
2249 (),
2250 ()
2251 );
2252 CALORECGPU_CMC_MOMENT_CALC
2253 (MZ,
2254 (),
2255 (CMCTemporaries::mZ(p.clusters_arr, cluster) = 0.f;
2256 CMCTemporaries::mZAux(p.clusters_arr, cluster) = 0.f;),
2257 (ToLoad::CellZ),
2258 (ToLoad::WeightedCellPositionNormalization),
2259 (const float mz = data.w_E_r_dir * data.z;
2260 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::mZ_ptr(p.clusters_arr, cluster),
2261 CMCTemporaries::mZAux_ptr(p.clusters_arr, cluster),
2262 mz);
2263 ),
2264 (),
2265 ()
2266 );
2267
2268
2269 CALORECGPU_CMC_MOMENT_CALC
2270 (NumPositiveEnergyCells,
2271 (),
2272 (CMCTemporaries::numPositiveEnergyCells(p.clusters_arr, cluster) = 0;),
2273 (),
2274 (ToLoad::WeightedEnergyOrNegative),
2275 (if (data.weighted_energy_or_negative > 0)
2276 {
2277 atomicAdd(&(CMCTemporaries::numPositiveEnergyCells(p.clusters_arr, cluster)), 1);
2278 }
2279 ),
2280 (),
2281 ()
2282 );
2283
2284 CALORECGPU_CMC_MOMENT_CALC
2285 (SeedCellPhi,
2286 (ToLoad::SeedCellGeometryPhi),
2287 (CMCTemporaries::seedCellPhi(p.clusters_arr, cluster) = data.seed_cell_phi_coordinate;),
2288 (),
2289 (),
2290 (),
2291 (),
2292 ()
2293 );
2294
2295 CALORECGPU_CMC_MOMENT_CALC
2296 (SumAbsEnergyNonMoments,
2297 (),
2298 (CMCTemporaries::sumAbsEnergyNonMoments(p.clusters_arr, cluster) = 0.f;
2299 CMCTemporaries::sumAbsEnergyNonMomentsAux(p.clusters_arr, cluster) = 0.f;
2300 ),
2301 (ToLoad::CellAbsEnergy),
2302 (),
2303 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::sumAbsEnergyNonMoments_ptr(p.clusters_arr, cluster),
2304 CMCTemporaries::sumAbsEnergyNonMomentsAux_ptr(p.clusters_arr, cluster),
2305 data.abs_energy * data.weight);
2306 ),
2307 (),
2308 ()
2309 );
2310
2311 CALORECGPU_CMC_MOMENT_CALC
2312 (SumSquareEnergies,
2313 (),
2314 (CMCTemporaries::sumSquareEnergies(p.clusters_arr, cluster) = 0.f;
2315 CMCTemporaries::sumSquareEnergiesAux(p.clusters_arr, cluster) = 0.f;
2316 ),
2317 (),
2318 (ToLoad::SquareWeightedEnergy),
2319 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::sumSquareEnergies_ptr(p.clusters_arr, cluster),
2320 CMCTemporaries::sumSquareEnergiesAux_ptr(p.clusters_arr, cluster),
2321 data.square_w_E);
2322 ),
2323 (),
2324 ()
2325 );
2326
2327 CALORECGPU_CMC_MOMENT_CALC
2328 (TimeNormalization,
2329 (),
2330 (CMCTemporaries::timeNormalization(p.clusters_arr, cluster) = 0.f;
2331 CMCTemporaries::timeNormalizationAux(p.clusters_arr, cluster) = 0.f;
2332 ),
2333 (ToLoad::CellTimeMomentsCheck),
2334 (ToLoad::SquaredWeightedNonMomentsEnergy),
2335 (if (data.time_moments_check)
2336 {
2337 CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::timeNormalization_ptr(p.clusters_arr, cluster),
2338 CMCTemporaries::timeNormalizationAux_ptr(p.clusters_arr, cluster),
2339 data.squared_normE);
2340 }
2341 ),
2342 (),
2343 ()
2344 );
2345
2346
2347 //+---------------------------------------------------+
2348 //| Per-Sampling Temporaries: |
2349 //+---------------------------------------------------+
2350
2351 template <int num, int delta = 0>
2352 CALORECGPU_CMC_MOMENT_CALC
2353 (AbsoluteEnergyPerSampleSeveral,
2354 (ToLoad::SamplingFromMomentIndex),
2355 (const int offset = data.sampling * num + delta;
2356 for (int i = 0; i < num; ++i)
2357 {
2358 CMCTemporaries::absoluteEnergyPerSample(p.clusters_arr, offset + i, cluster)
2359 = 0.f;
2360 CMCTemporaries::absoluteEnergyPerSampleAux(p.clusters_arr, offset + i, cluster) = 0.f;
2361 }
2362 ),
2363 (ToLoad::CellSampling, ToLoad::CellAbsEnergy),
2364 (),
2365 (CaloRecGPU::Helpers::device_kahan_babushka_neumaier(CMCTemporaries::absoluteEnergyPerSample_ptr(p.clusters_arr, data.sampling, cluster),
2366 CMCTemporaries::absoluteEnergyPerSampleAux_ptr(p.clusters_arr, data.sampling, cluster),
2367 data.abs_energy * data.weight);
2368 ),
2369 (),
2370 ()
2371 );
2372
2373 using AbsoluteEnergyPerSample = AbsoluteEnergyPerSampleSeveral<1, 0>;
2374
2375 template <int num, int delta = 0>
2376 CALORECGPU_CMC_MOMENT_CALC
2377 (MaxEnergyAndCellPerSampleSeveral,
2378 (ToLoad::SamplingFromMomentIndex),
2379 (const int offset = data.sampling * num + delta;
2380 for (int i = 0; i < num; ++i)
2381 {
2382 CMCTemporaries::maxEnergyAndCellPerSample(p.clusters_arr, offset + i, cluster)
2383 = 0ULL;
2384 }
2385 ),
2386 (ToLoad::CellSampling, ToLoad::CellEnergy),
2387 (),
2388 (const unsigned int energy_pattern = __float_as_uint(data.energy * data.weight);
2389 unsigned long long int E_and_cell = FloatingPointHelpers::StandardFloat::to_total_ordering(energy_pattern);
2390 E_and_cell = (E_and_cell << 32) | cell;
2391 atomicMax(&(CMCTemporaries::maxEnergyAndCellPerSample(p.clusters_arr, data.sampling, cluster)), E_and_cell);
2392 ),
2393 (),
2394 ()
2395 );
2396
2397 using MaxEnergyAndCellPerSample = MaxEnergyAndCellPerSampleSeveral<1, 0>;
2398
2399 template <int num, int delta = 0>
2400 CALORECGPU_CMC_MOMENT_CALC
2401 (MaxECellPerSampleSeveral,
2402 (ToLoad::SamplingFromMomentIndex/*, ToLoad::MaxEnergyAndCellPerSample*/),
2403 (constexpr unsigned long long int total_ordering_zero = FloatingPointHelpers::StandardFloat::to_total_ordering(0x00000000U);
2404 //Floating point 0 is all 0.
2405 constexpr unsigned long long int comparison = (total_ordering_zero << 32) | 0xFFFFFFFFU;
2406 const int offset = data.sampling * num + delta;
2407 for (int i = 0; i < num; ++i)
2408 {
2409 const unsigned long long max_energy_and_cell = CMCTemporaries::maxEnergyAndCellPerSample(p.clusters_arr, offset + i, cluster);
2410 const int cell = (max_energy_and_cell > comparison ? ((int) (max_energy_and_cell & 0x7FFFFFFFU)) : -1);
2411 CMCTemporaries::maxECellPerSample(p.clusters_arr, offset + i, cluster) = cell;
2412 }
2413 ),
2414 (),
2415 (),
2416 (),
2417 (),
2418 ()
2419 );
2420
2421 using MaxECellPerSample = MaxECellPerSampleSeveral<1, 0>;
2422 }
2423}
2424
2425//Possible TO-DO: Write a system that would validate the dependencies
2426// so we could have an extra layer of protection
2427// if/when reordering the operations.
2428
2429#endif
eta
Scalar eta() const
pseudorapidity method
Definition AmgMatrixBasePlugin.h:83
deltaPhi
Scalar deltaPhi(const MatrixBase< Derived > &vec) const
Definition AmgMatrixBasePlugin.h:112
phi
Scalar phi() const
phi method
Definition AmgMatrixBasePlugin.h:67
FPHelpers.h
Contains functions to deal with arbitrary IEEE754-like floating point formats.
CALORECGPU_CMC_LOAD_SIMPLE_CLUSTER_INFO
#define CALORECGPU_CMC_LOAD_SIMPLE_CLUSTER_INFO(NAME, VARNAME, PROPNAME)
Definition GPUClusterInfoAndMomentsCalculatorImplHelper.h:696
CALORECGPU_CMC_LOAD_SIMPLE_CELL_INFO
#define CALORECGPU_CMC_LOAD_SIMPLE_CELL_INFO(NAME, VARNAME, PROPNAME)
Definition GPUClusterInfoAndMomentsCalculatorImplHelper.h:680
CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO
#define CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(NAME, VARNAME, PROPNAME)
Definition GPUClusterInfoAndMomentsCalculatorImplHelper.h:722
CALORECGPU_CMC_LOAD_SIMPLE_MOMENT_INFO
#define CALORECGPU_CMC_LOAD_SIMPLE_MOMENT_INFO(NAME, VARNAME, PROPNAME)
Definition GPUClusterInfoAndMomentsCalculatorImplHelper.h:704
CALORECGPU_CMC_MOMENT_CALC
#define CALORECGPU_CMC_MOMENT_CALC(NAME, BEFORELOAD, BEFOREEXEC, CELLLOAD, CLUSTERLOAD, CELLEXEC, AFTERLOAD, AFTEREXEC)
Definition GPUClusterInfoAndMomentsCalculatorImplHelper.h:1201
CALORECGPU_CMC_LOAD_SIMPLE_PER_SAMPLING_TEMPORARY_INFO
#define CALORECGPU_CMC_LOAD_SIMPLE_PER_SAMPLING_TEMPORARY_INFO(NAME, VARNAME, PROPNAME)
Definition GPUClusterInfoAndMomentsCalculatorImplHelper.h:732
CALORECGPU_CMC_LOAD_SIMPLE_GEOMETRY_INFO
#define CALORECGPU_CMC_LOAD_SIMPLE_GEOMETRY_INFO(NAME, VARNAME, PROPNAME)
Definition GPUClusterInfoAndMomentsCalculatorImplHelper.h:688
CALORECGPU_CMC_LOAD
#define CALORECGPU_CMC_LOAD(NAME, NEEDED, PREVNEEDED, VARS, INIT)
Definition GPUClusterInfoAndMomentsCalculatorImplHelper.h:669
data
char data[hepevt_bytes_allocation_ATLAS]
Definition HepEvt.cxx:11
if
if(febId1==febId2)
Definition LArRodBlockPhysicsV0.cxx:567
operator+
NavigableIterator< CONT, RPAR, COLL > operator+(typename NavigableIterator< CONT, RPAR, COLL >::difference_type n, NavigableIterator< CONT, RPAR, COLL > a)
Definition NavigableIterator.h:219
CALORECGPU_TEMP2DARR_2
#define CALORECGPU_TEMP2DARR_2(TEMPNAME, BASEVAR1, BASEVAR2, TYPE)
Defines a temporary variable of type TYPE with name TEMPNAME that will will be split across two 2D ar...
Definition TemporaryHelpers.h:341
CALORECGPU_TEMPARR_1
#define CALORECGPU_TEMPARR_1(TEMPNAME, BASEVAR, TYPE)
Defines a temporary variable of type TYPE with name TEMPNAME that will be stored in the 1D array acce...
Definition TemporaryHelpers.h:269
CALORECGPU_TEMP2DARR_1
#define CALORECGPU_TEMP2DARR_1(TEMPNAME, BASEVAR, TYPE)
Defines a temporary variable of type TYPE with name TEMPNAME that will be stored in the 2D array (i.
Definition TemporaryHelpers.h:319
CALORECGPU_TEMPARR_2
#define CALORECGPU_TEMPARR_2(TEMPNAME, BASEVAR1, BASEVAR2, TYPE)
Defines a temporary variable of type TYPE with name TEMPNAME that will be split across two 1D arrays ...
Definition TemporaryHelpers.h:291
y
#define y
x
#define x
z
#define z
operator^
constexpr auto operator^(E lhs, F rhs)
operator^
Definition bitmask.h:173
operator|
constexpr auto operator|(E lhs, F rhs)
operator|
Definition bitmask.h:160
max
#define max(a, b)
Definition cfImp.cxx:41
r
int r
Definition globals.cxx:22
count
int count(std::string s, const std::string &regx)
count how many occurances of a regx are in a string
Definition hcg.cxx:146
CaloRecGPU::Helpers::regularize_angle
static CUDA_HOS_DEV float regularize_angle(const float b, const float a=0.f)
Definition Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:421
CaloRecGPU::Helpers::corrected_magn_cross_product
static CUDA_HOS_DEV float corrected_magn_cross_product(const float a1, const float a2, const float a3, const float b1, const float b2, const float b3)
Definition Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:605
CaloRecGPU::Helpers::product_sum_cornea_harrison_tang
static CUDA_HOS_DEV float product_sum_cornea_harrison_tang(const float a, const float b, const float c, const float d)
Definition Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:549
CaloRecGPU::Helpers::corrected_dot_product
static CUDA_HOS_DEV float corrected_dot_product(const float a_1, const float a_2, const float a_3, const float b_1, const float b_2, const float b_3)
Definition Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:564
CaloRecGPU::Helpers::sum_kahan_babushka_neumaier
CUDA_HOS_DEV float sum_kahan_babushka_neumaier(const Floats &... fs)
Adds a list of floats using the Kahan-Babushka-Neumaier algorithm for greater precision (at the cost ...
Definition Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:511
CaloRecGPU::Helpers::CUDA_kernel_object
SimpleHolder< T, MemoryContext::CUDAGPU, false > CUDA_kernel_object
Non-owning pointer to an object of type T in CUDA GPU memory.
Definition Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1667
CaloRecGPU::NCaloCells
constexpr int NCaloCells
Definition BaseDefinitions.h:12
FlavorTagInference::combine
size_t combine(size_t lhs, size_t rhs)
Definition hash.h:21
LArNewCalib_DelayDump_OFC_Cali.idx
idx
Definition LArNewCalib_DelayDump_OFC_Cali.py:69
ParticleGun_FastCalo_ChargeFlip_Config.energy
energy
Definition ParticleGun_FastCalo_ChargeFlip_Config.py:78
TSU::T
unsigned long long T
Definition L1TopoDataTypes.h:35
dqt_zlumi_pandas.weight
int weight
Definition dqt_zlumi_pandas.py:190
hist_file_dump.f
f
Definition hist_file_dump.py:140
trigbs_pickEvents.num
int num
Definition trigbs_pickEvents.py:76
CaloRecGPU::QualityProvenance
Just two uint16_t bit-packed onto a uint32_t.
Definition EventInfoDefinitions.h:116
FloatingPointHelpers::IEEE754_like< 23, 8 >::to_total_ordering
static constexpr T to_total_ordering(const T pattern)
Definition FPHelpers.h:444
number
std::string number(const double &d, const std::string &s)
Definition utils.cxx:186
Generated on for ATLAS Offline Software by doxygen 1.14.0