Warning, /RecoLocalCalo/EcalRecProducers/plugins/KernelHelpers.cu is written in an unsupported language. File is not indexed.
0001 #include "DataFormats/EcalDetId/interface/EBDetId.h"
0002 #include "DataFormats/EcalDetId/interface/EEDetId.h"
0003
0004 #include "KernelHelpers.h"
0005
0006 namespace ecal {
0007 namespace reconstruction {
0008
0009 namespace internal {
0010
0011 namespace barrel {
0012
0013 __device__ __forceinline__ bool positiveZ(uint32_t id) { return id & 0x10000; }
0014
0015 __device__ __forceinline__ uint32_t ietaAbs(uint32_t id) { return (id >> 9) & 0x7F; }
0016
0017 __device__ __forceinline__ uint32_t iphi(uint32_t id) { return id & 0x1FF; }
0018
0019 __device__ int dccFromSm(int ism) {
0020 int iz = 1;
0021 if (ism > 18)
0022 iz = -1;
0023 if (iz == -1)
0024 ism -= 18;
0025 int idcc = 9 + ism;
0026 if (iz == +1)
0027 idcc += 18;
0028 return idcc;
0029 }
0030
0031 __device__ int sm(int ieta, int iphi) {
0032 int iz = 1;
0033 if (ieta < 0)
0034 iz = -1;
0035 ieta *= iz;
0036 int iphi_ = iphi;
0037 if (iphi_ > 360)
0038 iphi_ -= 360;
0039 int ism = (iphi_ - 1) / 20 + 1;
0040 if (iz == -1)
0041 ism += 18;
0042 return ism;
0043 }
0044
0045 __device__ int dcc(int ieta, int iphi) {
0046 int ism = sm(ieta, iphi);
0047 return dccFromSm(ism);
0048 }
0049
0050 __device__ int lm_channel(int iX, int iY) {
0051 static const int idx_[] = {
0052 // clang-format off
0053 // 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16
0054 1, 2, 2, 2, 2, 4, 4, 4, 4, 6, 6, 6, 6, 8, 8, 8, 8, // 3
0055 1, 2, 2, 2, 2, 4, 4, 4, 4, 6, 6, 6, 6, 8, 8, 8, 8, // 2
0056 1, 3, 3, 3, 3, 5, 5, 5, 5, 7, 7, 7, 7, 9, 9, 9, 9, // 1
0057 1, 3, 3, 3, 3, 5, 5, 5, 5, 7, 7, 7, 7, 9, 9, 9, 9 // 0
0058 // 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16
0059 // clang-format on
0060 };
0061
0062 int il, ic, ii;
0063 const int iym = 4;
0064 const int ixm = 17;
0065 int iX_ = iX + 1;
0066 int iY_ = iY + 1;
0067 il = iym - iY_;
0068 ic = iX_ - 1;
0069 ii = il * ixm + ic;
0070 if (ii < 0 || ii > (int)(sizeof(idx_) / sizeof(int))) {
0071 return -1;
0072 };
0073 return idx_[ii];
0074 }
0075
0076 __device__ int localCoord_x(int ieta, int iphi) {
0077 int iz = 1;
0078 if (ieta < 0) {
0079 iz = -1;
0080 }
0081 ieta *= iz;
0082 int ix = ieta - 1;
0083
0084 return ix;
0085 }
0086
0087 __device__ int localCoord_y(int ieta, int iphi) {
0088 int iz = 1;
0089 if (ieta < 0) {
0090 iz = -1;
0091 }
0092 int iphi_ = iphi;
0093 if (iphi_ > 360) {
0094 iphi_ -= 360;
0095 }
0096 int iy = (iphi_ - 1) % 20;
0097 if (iz == -1) {
0098 iy = 19 - iy;
0099 }
0100
0101 return iy;
0102 }
0103
0104 __device__ int lmmod(int ieta, int iphi) {
0105 int ix = localCoord_x(ieta, iphi);
0106 int iy = localCoord_y(ieta, iphi);
0107
0108 return lm_channel(ix / 5, iy / 5);
0109 }
0110
0111 __device__ int side(int ieta, int iphi) {
0112 int ilmmod = lmmod(ieta, iphi);
0113 return (ilmmod % 2 == 0) ? 1 : 0;
0114 }
0115
0116 } // namespace barrel
0117
0118 } // namespace internal
0119
0120 __device__ uint32_t hashedIndexEB(uint32_t id) {
0121 using namespace internal::barrel;
0122 return (EBDetId::MAX_IETA + (positiveZ(id) ? ietaAbs(id) - 1 : -ietaAbs(id))) * EBDetId::MAX_IPHI + iphi(id) - 1;
0123 }
0124
0125 //
0126 // https://cmssdt.cern.ch/lxr/source/CalibCalorimetry/EcalLaserAnalyzer/src/MEEBGeom.cc
0127 // function: "lmr"
0128
0129 __device__ int laser_monitoring_region_EB(uint32_t id) {
0130 using namespace internal::barrel;
0131
0132 int ieta;
0133 if (positiveZ(id)) {
0134 ieta = ietaAbs(id);
0135 } else {
0136 ieta = -ietaAbs(id);
0137 }
0138
0139 int idcc = dcc(ieta, (int)(iphi(id)));
0140 int ism = idcc - 9;
0141
0142 int iside = side(ieta, (int)(iphi(id)));
0143
0144 return (1 + 2 * (ism - 1) + iside);
0145 }
0146
0147 namespace internal {
0148
0149 namespace endcap {
0150
0151 __device__ __forceinline__ uint32_t ix(uint32_t id) { return (id >> 7) & 0x7F; }
0152
0153 __device__ __forceinline__ uint32_t iy(uint32_t id) { return id & 0x7F; }
0154
0155 __device__ __forceinline__ bool positiveZ(uint32_t id) { return id & 0x4000; }
0156
0157 // these constants come from EE Det Id
0158 __constant__ const unsigned short kxf[] = {
0159 41, 51, 41, 51, 41, 51, 36, 51, 36, 51, 26, 51, 26, 51, 26, 51, 21, 51, 21, 51, 21, 51, 21, 51, 21,
0160 51, 16, 51, 16, 51, 14, 51, 14, 51, 14, 51, 14, 51, 14, 51, 9, 51, 9, 51, 9, 51, 9, 51, 9, 51,
0161 6, 51, 6, 51, 6, 51, 6, 51, 6, 51, 6, 51, 6, 51, 6, 51, 6, 51, 6, 51, 4, 51, 4, 51, 4,
0162 51, 4, 51, 4, 56, 1, 58, 1, 59, 1, 60, 1, 61, 1, 61, 1, 62, 1, 62, 1, 62, 1, 62, 1, 62,
0163 1, 62, 1, 62, 1, 62, 1, 62, 1, 62, 1, 61, 1, 61, 1, 60, 1, 59, 1, 58, 4, 56, 4, 51, 4,
0164 51, 4, 51, 4, 51, 6, 51, 6, 51, 6, 51, 6, 51, 6, 51, 6, 51, 6, 51, 6, 51, 6, 51, 6, 51,
0165 9, 51, 9, 51, 9, 51, 9, 51, 9, 51, 14, 51, 14, 51, 14, 51, 14, 51, 14, 51, 16, 51, 16, 51, 21,
0166 51, 21, 51, 21, 51, 21, 51, 21, 51, 26, 51, 26, 51, 26, 51, 36, 51, 36, 51, 41, 51, 41, 51, 41, 51};
0167
0168 __constant__ const unsigned short kdi[] = {
0169 0, 10, 20, 30, 40, 50, 60, 75, 90, 105, 120, 145, 170, 195, 220, 245, 270,
0170 300, 330, 360, 390, 420, 450, 480, 510, 540, 570, 605, 640, 675, 710, 747, 784, 821,
0171 858, 895, 932, 969, 1006, 1043, 1080, 1122, 1164, 1206, 1248, 1290, 1332, 1374, 1416, 1458, 1500,
0172 1545, 1590, 1635, 1680, 1725, 1770, 1815, 1860, 1905, 1950, 1995, 2040, 2085, 2130, 2175, 2220, 2265,
0173 2310, 2355, 2400, 2447, 2494, 2541, 2588, 2635, 2682, 2729, 2776, 2818, 2860, 2903, 2946, 2988, 3030,
0174 3071, 3112, 3152, 3192, 3232, 3272, 3311, 3350, 3389, 3428, 3467, 3506, 3545, 3584, 3623, 3662, 3701,
0175 3740, 3779, 3818, 3857, 3896, 3935, 3974, 4013, 4052, 4092, 4132, 4172, 4212, 4253, 4294, 4336, 4378,
0176 4421, 4464, 4506, 4548, 4595, 4642, 4689, 4736, 4783, 4830, 4877, 4924, 4969, 5014, 5059, 5104, 5149,
0177 5194, 5239, 5284, 5329, 5374, 5419, 5464, 5509, 5554, 5599, 5644, 5689, 5734, 5779, 5824, 5866, 5908,
0178 5950, 5992, 6034, 6076, 6118, 6160, 6202, 6244, 6281, 6318, 6355, 6392, 6429, 6466, 6503, 6540, 6577,
0179 6614, 6649, 6684, 6719, 6754, 6784, 6814, 6844, 6874, 6904, 6934, 6964, 6994, 7024, 7054, 7079, 7104,
0180 7129, 7154, 7179, 7204, 7219, 7234, 7249, 7264, 7274, 7284, 7294, 7304, 7314};
0181
0182 __device__ int quadrant(int iX, int iY) {
0183 bool near = iX >= 11;
0184 bool far = !near;
0185 bool top = iY >= 11;
0186 bool bot = !top;
0187
0188 int iquad = 0;
0189 if (near && top)
0190 iquad = 1;
0191 if (far && top)
0192 iquad = 2;
0193 if (far && bot)
0194 iquad = 3;
0195 if (near && bot)
0196 iquad = 4;
0197
0198 return iquad;
0199 }
0200
0201 __device__ int sector(int iX, int iY) {
0202 // Y (towards the surface)
0203 // T
0204 // |
0205 // |
0206 // |
0207 // o---------| X (towards center of LHC)
0208 //
0209 static const int idx_[] = {
0210 // clang-format off
0211 // 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20
0212 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 9, 9, 9, 0, 0, 0, 0, 0, 0, 0, // 20
0213 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 9, 9, 9, 9, 9, 9, 0, 0, 0, 0, // 19
0214 0, 0, 0, 2, 1, 1, 1, 1, 1, 1, 9, 9, 9, 9, 9, 9, 8, 0, 0, 0, // 18
0215 0, 0, 2, 2, 2, 1, 1, 1, 1, 1, 9, 9, 9, 9, 9, 8, 8, 8, 0, 0, // 17
0216 0, 2, 2, 2, 2, 1, 1, 1, 1, 1, 9, 9, 9, 9, 9, 8, 8, 8, 8, 0, // 16
0217 0, 2, 2, 2, 2, 2, 1, 1, 1, 1, 9, 9, 9, 9, 8, 8, 8, 8, 8, 0, // 15
0218 0, 2, 2, 2, 2, 2, 2, 1, 1, 1, 9, 9, 9, 8, 8, 8, 8, 8, 8, 0, // 14
0219 2, 2, 2, 2, 2, 2, 2, 2, 1, 1, 9, 9, 8, 8, 8, 8, 8, 8, 8, 8, // 13
0220 3, 3, 2, 2, 2, 2, 2, 2, 2, 0, 0, 8, 8, 8, 8, 8, 8, 8, 7, 7, // 12
0221 3, 3, 3, 3, 3, 3, 3, 2, 0, 0, 0, 0, 8, 7, 7, 7, 7, 7, 7, 7, // 11
0222 3, 3, 3, 3, 3, 3, 3, 3, 0, 0, 0, 0, 7, 7, 7, 7, 7, 7, 7, 7, // 10
0223 3, 3, 3, 3, 3, 3, 3, 4, 4, 0, 0, 6, 6, 7, 7, 7, 7, 7, 7, 7, // 9
0224 3, 3, 3, 3, 3, 3, 4, 4, 4, 5, 5, 6, 6, 6, 7, 7, 7, 7, 7, 7, // 8
0225 0, 3, 3, 3, 4, 4, 4, 4, 4, 5, 5, 6, 6, 6, 6, 6, 7, 7, 7, 0, // 7
0226 0, 3, 4, 4, 4, 4, 4, 4, 5, 5, 5, 5, 6, 6, 6, 6, 6, 6, 7, 0, // 6
0227 0, 4, 4, 4, 4, 4, 4, 4, 5, 5, 5, 5, 6, 6, 6, 6, 6, 6, 6, 0, // 5
0228 0, 0, 4, 4, 4, 4, 4, 4, 5, 5, 5, 5, 6, 6, 6, 6, 6, 6, 0, 0, // 4
0229 0, 0, 0, 4, 4, 4, 4, 5, 5, 5, 5, 5, 5, 6, 6, 6, 6, 0, 0, 0, // 3
0230 0, 0, 0, 0, 4, 4, 4, 5, 5, 5, 5, 5, 5, 6, 6, 6, 0, 0, 0, 0, // 2
0231 0, 0, 0, 0, 0, 0, 0, 5, 5, 5, 5, 5, 5, 0, 0, 0, 0, 0, 0, 0 // 1
0232 // 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20
0233 // clang-format on
0234 };
0235
0236 int iym, ixm, il, ic, ii;
0237 iym = 20;
0238 ixm = 20;
0239 int iX_ = iX;
0240 int iY_ = iY;
0241 il = iym - iY_;
0242 ic = iX_ - 1;
0243 ii = il * ixm + ic;
0244
0245 if (ii < 0 || ii > (int)(sizeof(idx_) / sizeof(int)) || idx_[ii] == 0) {
0246 return -1;
0247 };
0248 return idx_[ii];
0249 }
0250
0251 } // namespace endcap
0252
0253 } // namespace internal
0254
0255 __device__ uint32_t hashedIndexEE(uint32_t id) {
0256 using namespace internal::endcap;
0257
0258 const uint32_t jx(ix(id));
0259 const uint32_t jd(2 * (iy(id) - 1) + (jx - 1) / 50);
0260 return ((positiveZ(id) ? EEDetId::kEEhalf : 0) + kdi[jd] + jx - kxf[jd]);
0261 }
0262
0263 //
0264 // https://cmssdt.cern.ch/lxr/source/CalibCalorimetry/EcalLaserAnalyzer/src/MEEEGeom.cc
0265 // https://github.com/cms-sw/cmssw/blob/master/CalibCalorimetry/EcalLaserCorrection/src/EcalLaserDbService.cc
0266 //
0267
0268 __device__ int laser_monitoring_region_EE(uint32_t id) {
0269 using namespace internal::endcap;
0270
0271 // SuperCrysCoord
0272 uint32_t iX = (ix(id) - 1) / 5 + 1;
0273 uint32_t iY = (iy(id) - 1) / 5 + 1;
0274
0275 // Correct convention
0276 // * @param iz iz/zside index: -1 for EE-, +1 for EE+
0277 // https://github.com/cms-sw/cmssw/blob/master/DataFormats/EcalDetId/interface/EEDetId.h#L68-L71
0278 // zside in https://github.com/cms-sw/cmssw/blob/master/CalibCalorimetry/EcalLaserCorrection/src/EcalLaserDbService.cc#L63
0279 //
0280 int iz = positiveZ(id) ? 1 : -1;
0281
0282 int iquad = quadrant(iX, iY);
0283 int isect = sector(iX, iY);
0284 if (isect < 0)
0285 return -1;
0286
0287 int ilmr = 0;
0288 ilmr = isect - 6;
0289 if (ilmr <= 0)
0290 ilmr += 9;
0291 if (ilmr == 9)
0292 ilmr++;
0293 if (ilmr == 8 && iquad == 4)
0294 ilmr++;
0295 if (iz == +1)
0296 ilmr += 72;
0297 else
0298 ilmr += 82;
0299
0300 return ilmr;
0301 }
0302
0303 } // namespace reconstruction
0304 } // namespace ecal
