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