Back to home page

Project CMSSW displayed by LXR

 
 

    


File indexing completed on 2023-01-19 02:53:02

0001 #ifndef Geometry_CommonTopologies_SimplePixelTopology_h
0002 #define Geometry_CommonTopologies_SimplePixelTopology_h
0003 
0004 #include <array>
0005 #include <cstdint>
0006 #include <type_traits>
0007 #include "FWCore/Utilities/interface/HostDeviceConstant.h"
0008 
0009 namespace pixelTopology {
0010 
0011   constexpr auto maxNumberOfLadders = 160;
0012   constexpr uint32_t maxLayers = 28;
0013 
0014   template <typename TrackerTraits>
0015   struct AverageGeometryT {
0016     //
0017     float ladderZ[TrackerTraits::numberOfLaddersInBarrel];
0018     float ladderX[TrackerTraits::numberOfLaddersInBarrel];
0019     float ladderY[TrackerTraits::numberOfLaddersInBarrel];
0020     float ladderR[TrackerTraits::numberOfLaddersInBarrel];
0021     float ladderMinZ[TrackerTraits::numberOfLaddersInBarrel];
0022     float ladderMaxZ[TrackerTraits::numberOfLaddersInBarrel];
0023     float endCapZ[2];  // just for pos and neg Layer1
0024   };
0025 
0026   constexpr int16_t phi0p05 = 522;  // round(521.52189...) = phi2short(0.05);
0027   constexpr int16_t phi0p06 = 626;  // round(625.82270...) = phi2short(0.06);
0028   constexpr int16_t phi0p07 = 730;  // round(730.12648...) = phi2short(0.07);
0029   constexpr int16_t phi0p09 = 900;
0030 
0031   constexpr uint16_t last_barrel_layer = 3;  // this is common between all the topologies
0032 
0033   template <class Function, std::size_t... Indices>
0034   constexpr auto map_to_array_helper(Function f, std::index_sequence<Indices...>)
0035       -> std::array<std::invoke_result_t<Function, std::size_t>, sizeof...(Indices)> {
0036     return {{f(Indices)...}};
0037   }
0038 
0039   template <int N, class Function>
0040   constexpr auto map_to_array(Function f) -> std::array<std::invoke_result_t<Function, std::size_t>, N> {
0041     return map_to_array_helper(f, std::make_index_sequence<N>{});
0042   }
0043 
0044   template <typename TrackerTraits>
0045   constexpr uint16_t findMaxModuleStride() {
0046     bool go = true;
0047     int n = 2;
0048     while (go) {
0049       for (uint8_t i = 1; i < TrackerTraits::numberOfLayers + 1; ++i) {
0050         if (TrackerTraits::layerStart[i] % n != 0) {
0051           go = false;
0052           break;
0053         }
0054       }
0055       if (!go)
0056         break;
0057       n *= 2;
0058     }
0059     return n / 2;
0060   }
0061 
0062   template <typename TrackerTraits>
0063   constexpr uint16_t maxModuleStride = findMaxModuleStride<TrackerTraits>();
0064 
0065   template <typename TrackerTraits>
0066   constexpr uint8_t findLayer(uint32_t detId, uint8_t sl = 0) {
0067     for (uint8_t i = sl; i < TrackerTraits::numberOfLayers + 1; ++i)
0068       if (detId < TrackerTraits::layerStart[i + 1])
0069         return i;
0070     return TrackerTraits::numberOfLayers + 1;
0071   }
0072 
0073   template <typename TrackerTraits>
0074   constexpr uint8_t findLayerFromCompact(uint32_t detId) {
0075     detId *= maxModuleStride<TrackerTraits>;
0076     for (uint8_t i = 0; i < TrackerTraits::numberOfLayers + 1; ++i)
0077       if (detId < TrackerTraits::layerStart[i + 1])
0078         return i;
0079     return TrackerTraits::numberOfLayers + 1;
0080   }
0081 
0082   template <typename TrackerTraits>
0083   constexpr uint32_t layerIndexSize = TrackerTraits::numberOfModules / maxModuleStride<TrackerTraits>;
0084 
0085   template <typename TrackerTraits>
0086 #ifdef __CUDA_ARCH__
0087   __device__
0088 #endif
0089       constexpr std::array<uint8_t, layerIndexSize<TrackerTraits>>
0090           layer = map_to_array<layerIndexSize<TrackerTraits>>(findLayerFromCompact<TrackerTraits>);
0091 
0092   template <typename TrackerTraits>
0093   constexpr uint8_t getLayer(uint32_t detId) {
0094     return layer<TrackerTraits>[detId / maxModuleStride<TrackerTraits>];
0095   }
0096 
0097   template <typename TrackerTraits>
0098   constexpr bool validateLayerIndex() {
0099     bool res = true;
0100     for (auto i = 0U; i < TrackerTraits::numberOfModules; ++i) {
0101       auto j = i / maxModuleStride<TrackerTraits>;
0102       res &= (layer<TrackerTraits>[j] < TrackerTraits::numberOfLayers);
0103       res &= (i >= TrackerTraits::layerStart[layer<TrackerTraits>[j]]);
0104       res &= (i < TrackerTraits::layerStart[layer<TrackerTraits>[j] + 1]);
0105     }
0106     return res;
0107   }
0108 
0109   template <typename TrackerTraits>
0110 #ifdef __CUDA_ARCH__
0111   __device__
0112 #endif
0113       constexpr inline uint32_t
0114       layerStart(uint32_t i) {
0115     return TrackerTraits::layerStart[i];
0116   }
0117 
0118   constexpr inline uint16_t divu52(uint16_t n) {
0119     n = n >> 2;
0120     uint16_t q = (n >> 1) + (n >> 4);
0121     q = q + (q >> 4) + (q >> 5);
0122     q = q >> 3;
0123     uint16_t r = n - q * 13;
0124     return q + ((r + 3) >> 4);
0125   }
0126 }  // namespace pixelTopology
0127 
0128 namespace phase1PixelTopology {
0129 
0130   using pixelTopology::phi0p05;
0131   using pixelTopology::phi0p06;
0132   using pixelTopology::phi0p07;
0133 
0134   constexpr uint32_t numberOfLayers = 28;
0135   constexpr int nPairs = 13 + 2 + 4;
0136   constexpr uint16_t numberOfModules = 1856;
0137 
0138   constexpr uint32_t max_ladder_bpx0 = 12;
0139   constexpr uint32_t first_ladder_bpx0 = 0;
0140   constexpr float module_length_bpx0 = 6.7f;
0141   constexpr float module_tolerance_bpx0 = 0.4f;  // projection to cylinder is inaccurate on BPIX1
0142   constexpr uint32_t max_ladder_bpx4 = 64;
0143   constexpr uint32_t first_ladder_bpx4 = 84;
0144   constexpr float radius_even_ladder = 15.815f;
0145   constexpr float radius_odd_ladder = 16.146f;
0146   constexpr float module_length_bpx4 = 6.7f;
0147   constexpr float module_tolerance_bpx4 = 0.2f;
0148   constexpr float barrel_z_length = 26.f;
0149   constexpr float forward_z_begin = 32.f;
0150 
0151   HOST_DEVICE_CONSTANT uint8_t layerPairs[2 * nPairs] = {
0152       0, 1, 0, 4, 0, 7,              // BPIX1 (3)
0153       1, 2, 1, 4, 1, 7,              // BPIX2 (6)
0154       4, 5, 7, 8,                    // FPIX1 (8)
0155       2, 3, 2, 4, 2, 7, 5, 6, 8, 9,  // BPIX3 & FPIX2 (13)
0156       0, 2, 1, 3,                    // Jumping Barrel (15)
0157       0, 5, 0, 8,                    // Jumping Forward (BPIX1,FPIX2)
0158       4, 6, 7, 9                     // Jumping Forward (19)
0159   };
0160 
0161   HOST_DEVICE_CONSTANT int16_t phicuts[nPairs]{phi0p05,
0162                                                phi0p07,
0163                                                phi0p07,
0164                                                phi0p05,
0165                                                phi0p06,
0166                                                phi0p06,
0167                                                phi0p05,
0168                                                phi0p05,
0169                                                phi0p06,
0170                                                phi0p06,
0171                                                phi0p06,
0172                                                phi0p05,
0173                                                phi0p05,
0174                                                phi0p05,
0175                                                phi0p05,
0176                                                phi0p05,
0177                                                phi0p05,
0178                                                phi0p05,
0179                                                phi0p05};
0180   HOST_DEVICE_CONSTANT float minz[nPairs] = {
0181       -20., 0., -30., -22., 10., -30., -70., -70., -22., 15., -30, -70., -70., -20., -22., 0, -30., -70., -70.};
0182   HOST_DEVICE_CONSTANT float maxz[nPairs] = {
0183       20., 30., 0., 22., 30., -10., 70., 70., 22., 30., -15., 70., 70., 20., 22., 30., 0., 70., 70.};
0184   HOST_DEVICE_CONSTANT float maxr[nPairs] = {
0185       20., 9., 9., 20., 7., 7., 5., 5., 20., 6., 6., 5., 5., 20., 20., 9., 9., 9., 9.};
0186 
0187   static constexpr uint32_t layerStart[numberOfLayers + 1] = {0,
0188                                                               96,
0189                                                               320,
0190                                                               672,  // barrel
0191                                                               1184,
0192                                                               1296,
0193                                                               1408,  // positive endcap
0194                                                               1520,
0195                                                               1632,
0196                                                               1744,  // negative endcap
0197                                                               numberOfModules};
0198 }  // namespace phase1PixelTopology
0199 
0200 namespace phase2PixelTopology {
0201 
0202   using pixelTopology::phi0p05;
0203   using pixelTopology::phi0p06;
0204   using pixelTopology::phi0p07;
0205   using pixelTopology::phi0p09;
0206 
0207   constexpr uint32_t numberOfLayers = 28;
0208   constexpr int nPairs = 23 + 6 + 14 + 8 + 4;  // include far forward layer pairs
0209   constexpr uint16_t numberOfModules = 3892;
0210 
0211   HOST_DEVICE_CONSTANT uint8_t layerPairs[2 * nPairs] = {
0212 
0213       0,  1,  0,  4,  0,  16,  //BPIX1 (3)
0214       1,  2,  1,  4,  1,  16,  //BPIX2 (6)
0215       2,  3,  2,  4,  2,  16,  //BPIX3 & Forward (9)
0216 
0217       4,  5,  5,  6,  6,  7,  7,  8,  8,  9,  9,  10, 10, 11,  //POS (16)
0218       16, 17, 17, 18, 18, 19, 19, 20, 20, 21, 21, 22, 22, 23,  //NEG (23)
0219 
0220       0,  2,  0,  5,  0,  17, 0,  6,  0,  18,  // BPIX1 Jump (28)
0221       1,  3,  1,  5,  1,  17, 1,  6,  1,  18,  // BPIX2 Jump (33)
0222 
0223       11, 12, 12, 13, 13, 14, 14, 15,  //Late POS (37)
0224       23, 24, 24, 25, 25, 26, 26, 27,  //Late NEG (41)
0225 
0226       4,  6,  5,  7,  6,  8,  7,  9,  8,  10, 9,  11, 10, 12,  //POS Jump (48)
0227       16, 18, 17, 19, 18, 20, 19, 21, 20, 22, 21, 23, 22, 24,  //NEG Jump (55)
0228   };
0229   HOST_DEVICE_CONSTANT uint32_t layerStart[numberOfLayers + 1] = {0,
0230                                                                   108,
0231                                                                   324,
0232                                                                   504,  //Barrel
0233                                                                   756,
0234                                                                   864,
0235                                                                   972,
0236                                                                   1080,
0237                                                                   1188,
0238                                                                   1296,
0239                                                                   1404,
0240                                                                   1512,
0241                                                                   1620,
0242                                                                   1796,
0243                                                                   1972,
0244                                                                   2148,  //Fp
0245                                                                   2324,
0246                                                                   2432,
0247                                                                   2540,
0248                                                                   2648,
0249                                                                   2756,
0250                                                                   2864,
0251                                                                   2972,
0252                                                                   3080,
0253                                                                   3188,
0254                                                                   3364,
0255                                                                   3540,
0256                                                                   3716,  //Np
0257                                                                   numberOfModules};
0258 
0259   HOST_DEVICE_CONSTANT int16_t phicuts[nPairs]{
0260       phi0p05, phi0p05, phi0p05, phi0p06, phi0p07, phi0p07, phi0p06, phi0p07, phi0p07, phi0p05, phi0p05,
0261       phi0p05, phi0p05, phi0p05, phi0p05, phi0p05, phi0p05, phi0p05, phi0p05, phi0p05, phi0p05, phi0p05,
0262       phi0p05, phi0p05, phi0p05, phi0p05, phi0p05, phi0p05, phi0p05, phi0p07, phi0p07, phi0p07, phi0p07,
0263       phi0p07, phi0p07, phi0p07, phi0p07, phi0p07, phi0p07, phi0p07, phi0p07, phi0p07, phi0p07, phi0p07,
0264       phi0p07, phi0p07, phi0p07, phi0p05, phi0p05, phi0p05, phi0p05, phi0p05, phi0p05, phi0p05, phi0p05};
0265 
0266   HOST_DEVICE_CONSTANT float minz[nPairs] = {
0267       -16.0, 4.0,   -22.0, -17.0, 6.0,   -22.0, -18.0, 11.0,  -22.0,  23.0,   30.0,   39.0,   50.0,   65.0,
0268       82.0,  109.0, -28.0, -35.0, -44.0, -55.0, -70.0, -87.0, -113.0, -16.,   7.0,    -22.0,  11.0,   -22.0,
0269       -17.0, 9.0,   -22.0, 13.0,  -22.0, 137.0, 173.0, 199.0, 229.0,  -142.0, -177.0, -203.0, -233.0, 23.0,
0270       30.0,  39.0,  50.0,  65.0,  82.0,  109.0, -28.0, -35.0, -44.0,  -55.0,  -70.0,  -87.0,  -113.0};
0271 
0272   HOST_DEVICE_CONSTANT float maxz[nPairs] = {
0273 
0274       17.0, 22.0,  -4.0,  17.0,  22.0,  -6.0,  18.0,  22.0,  -11.0,  28.0,   35.0,   44.0,   55.0,   70.0,
0275       87.0, 113.0, -23.0, -30.0, -39.0, -50.0, -65.0, -82.0, -109.0, 17.0,   22.0,   -7.0,   22.0,   -10.0,
0276       17.0, 22.0,  -9.0,  22.0,  -13.0, 142.0, 177.0, 203.0, 233.0,  -137.0, -173.0, -199.0, -229.0, 28.0,
0277       35.0, 44.0,  55.0,  70.0,  87.0,  113.0, -23.0, -30.0, -39.0,  -50.0,  -65.0,  -82.0,  -109.0};
0278 
0279   HOST_DEVICE_CONSTANT float maxr[nPairs] = {5.0, 5.0, 5.0, 7.0, 8.0, 8.0,  7.0, 7.0, 7.0, 6.0, 6.0, 6.0, 6.0, 5.0,
0280                                              6.0, 5.0, 6.0, 6.0, 6.0, 6.0,  5.0, 6.0, 5.0, 5.0, 5.0, 5.0, 5.0, 5.0,
0281                                              5.0, 8.0, 8.0, 8.0, 8.0, 6.0,  5.0, 5.0, 5.0, 6.0, 5.0, 5.0, 5.0, 9.0,
0282                                              9.0, 9.0, 8.0, 8.0, 8.0, 11.0, 9.0, 9.0, 9.0, 8.0, 8.0, 8.0, 11.0};
0283 }  // namespace phase2PixelTopology
0284 
0285 namespace pixelTopology {
0286 
0287   struct Phase2 {
0288     // types
0289     using hindex_type = uint32_t;  // FIXME from siPixelRecHitsHeterogeneousProduct
0290     using tindex_type = uint32_t;  // for tuples
0291     using cindex_type = uint32_t;  // for cells
0292 
0293     static constexpr uint32_t maxCellNeighbors = 64;
0294     static constexpr uint32_t maxCellTracks = 302;
0295     static constexpr uint32_t maxHitsOnTrack = 15;
0296     static constexpr uint32_t maxHitsOnTrackForFullFit = 6;
0297     static constexpr uint32_t avgHitsPerTrack = 7;
0298     static constexpr uint32_t maxCellsPerHit = 256;
0299     static constexpr uint32_t avgTracksPerHit = 10;
0300     static constexpr uint32_t maxNumberOfTuples = 256 * 1024;
0301     //this is well above thanks to maxNumberOfTuples
0302     static constexpr uint32_t maxHitsForContainers = avgHitsPerTrack * maxNumberOfTuples;
0303     static constexpr uint32_t maxNumberOfDoublets = 5 * 512 * 1024;
0304     static constexpr uint32_t maxNumOfActiveDoublets = maxNumberOfDoublets / 8;
0305     static constexpr uint32_t maxNumberOfQuadruplets = maxNumberOfTuples;
0306     static constexpr uint32_t maxDepth = 12;
0307     static constexpr uint32_t numberOfLayers = 28;
0308 
0309     static constexpr uint32_t maxSizeCluster = 2047;
0310 
0311     static constexpr uint32_t getDoubletsFromHistoMaxBlockSize = 64;  // for both x and y
0312     static constexpr uint32_t getDoubletsFromHistoMinBlocksPerMP = 16;
0313 
0314     static constexpr uint16_t last_bpix1_detIndex = 108;
0315     static constexpr uint16_t last_bpix2_detIndex = 324;
0316     static constexpr uint16_t last_barrel_detIndex = 504;
0317 
0318     static constexpr uint32_t maxPixInModule = 6000;
0319 
0320     static constexpr float moduleLength = 4.345f;
0321     static constexpr float endcapCorrection = 0.0f;
0322 
0323     static constexpr float xerr_barrel_l1_def = 0.00035f;
0324     static constexpr float yerr_barrel_l1_def = 0.00125f;
0325     static constexpr float xerr_barrel_ln_def = 0.00035f;
0326     static constexpr float yerr_barrel_ln_def = 0.00125f;
0327     static constexpr float xerr_endcap_def = 0.00060f;
0328     static constexpr float yerr_endcap_def = 0.00180f;
0329 
0330     static constexpr float bigPixXCorrection = 0.0f;
0331     static constexpr float bigPixYCorrection = 0.0f;
0332 
0333     static constexpr float dzdrFact = 8 * 0.0285 / 0.015;  // from dz/dr to "DY"
0334     static constexpr float z0Cut = 7.5f;
0335     static constexpr float doubletHardPt = 0.8f;
0336 
0337     static constexpr int minYsizeB1 = 25;
0338     static constexpr int minYsizeB2 = 15;
0339 
0340     static constexpr int nPairsMinimal = 33;
0341     static constexpr int nPairsFarForwards = nPairsMinimal + 8;  // include barrel "jumping" layer pairs
0342     static constexpr int nPairs = phase2PixelTopology::nPairs;   // include far forward layer pairs
0343 
0344     static constexpr int maxDYsize12 = 12;
0345     static constexpr int maxDYsize = 10;
0346     static constexpr int maxDYPred = 20;
0347 
0348     static constexpr uint16_t numberOfModules = 3892;
0349 
0350     static constexpr uint16_t clusterBinning = 1024;
0351     static constexpr uint16_t clusterBits = 10;
0352 
0353     static constexpr uint16_t numberOfModulesInBarrel = 756;
0354     static constexpr uint16_t numberOfModulesInLadder = 9;
0355     static constexpr uint16_t numberOfLaddersInBarrel = numberOfModulesInBarrel / numberOfModulesInLadder;
0356 
0357     static constexpr uint16_t firstEndcapPos = 4;
0358     static constexpr uint16_t firstEndcapNeg = 16;
0359 
0360     static constexpr int16_t xOffset = -1e4;  //not used actually, to suppress static analyzer warnings
0361 
0362     static constexpr char const *nameModifier = "Phase2";
0363 
0364     static constexpr uint32_t const *layerStart = phase2PixelTopology::layerStart;
0365     static constexpr float const *minz = phase2PixelTopology::minz;
0366     static constexpr float const *maxz = phase2PixelTopology::maxz;
0367     static constexpr float const *maxr = phase2PixelTopology::maxr;
0368 
0369     static constexpr uint8_t const *layerPairs = phase2PixelTopology::layerPairs;
0370     static constexpr int16_t const *phicuts = phase2PixelTopology::phicuts;
0371 
0372     static constexpr inline bool isBigPixX(uint16_t px) { return false; }
0373     static constexpr inline bool isBigPixY(uint16_t py) { return false; }
0374 
0375     static constexpr inline uint16_t localX(uint16_t px) { return px; }
0376     static constexpr inline uint16_t localY(uint16_t py) { return py; }
0377   };
0378 
0379   struct Phase1 {
0380     // types
0381     using hindex_type = uint32_t;  // FIXME from siPixelRecHitsHeterogeneousProduct
0382     using tindex_type = uint16_t;  // for tuples
0383     using cindex_type = uint32_t;  // for cells
0384 
0385     static constexpr uint32_t maxCellNeighbors = 36;
0386     static constexpr uint32_t maxCellTracks = 48;
0387     static constexpr uint32_t maxHitsOnTrack = 10;
0388     static constexpr uint32_t maxHitsOnTrackForFullFit = 6;
0389     static constexpr uint32_t avgHitsPerTrack = 5;
0390     static constexpr uint32_t maxCellsPerHit = 256;
0391     static constexpr uint32_t avgTracksPerHit = 6;
0392     static constexpr uint32_t maxNumberOfTuples = 32 * 1024;
0393     static constexpr uint32_t maxHitsForContainers = avgHitsPerTrack * maxNumberOfTuples;
0394     static constexpr uint32_t maxNumberOfDoublets = 512 * 1024;
0395     static constexpr uint32_t maxNumOfActiveDoublets = maxNumberOfDoublets / 8;
0396     static constexpr uint32_t maxNumberOfQuadruplets = maxNumberOfTuples;
0397     static constexpr uint32_t maxDepth = 6;
0398     static constexpr uint32_t numberOfLayers = 10;
0399 
0400     static constexpr uint32_t maxSizeCluster = 1023;
0401 
0402     static constexpr uint32_t getDoubletsFromHistoMaxBlockSize = 64;  // for both x and y
0403     static constexpr uint32_t getDoubletsFromHistoMinBlocksPerMP = 16;
0404 
0405     static constexpr uint16_t last_bpix1_detIndex = 96;
0406     static constexpr uint16_t last_bpix2_detIndex = 320;
0407     static constexpr uint16_t last_barrel_detIndex = 1184;
0408 
0409     static constexpr uint32_t maxPixInModule = 6000;
0410 
0411     static constexpr float moduleLength = 6.7f;
0412     static constexpr float endcapCorrection = 1.5f;
0413 
0414     static constexpr float xerr_barrel_l1_def = 0.00200f;
0415     static constexpr float yerr_barrel_l1_def = 0.00210f;
0416     static constexpr float xerr_barrel_ln_def = 0.00200f;
0417     static constexpr float yerr_barrel_ln_def = 0.00210f;
0418     static constexpr float xerr_endcap_def = 0.0020f;
0419     static constexpr float yerr_endcap_def = 0.00210f;
0420 
0421     static constexpr float bigPixXCorrection = 1.0f;
0422     static constexpr float bigPixYCorrection = 8.0f;
0423 
0424     static constexpr float dzdrFact = 8 * 0.0285 / 0.015;  // from dz/dr to "DY"
0425     static constexpr float z0Cut = 12.f;
0426     static constexpr float doubletHardPt = 0.5f;
0427 
0428     static constexpr int minYsizeB1 = 36;
0429     static constexpr int minYsizeB2 = 28;
0430 
0431     static constexpr int nPairsForQuadruplets = 13;                     // quadruplets require hits in all layers
0432     static constexpr int nPairsForTriplets = nPairsForQuadruplets + 2;  // include barrel "jumping" layer pairs
0433     static constexpr int nPairs = nPairsForTriplets + 4;                // include forward "jumping" layer pairs
0434 
0435     static constexpr int maxDYsize12 = 28;
0436     static constexpr int maxDYsize = 20;
0437     static constexpr int maxDYPred = 20;
0438 
0439     static constexpr uint16_t numberOfModules = 1856;
0440 
0441     static constexpr uint16_t numRowsInRoc = 80;
0442     static constexpr uint16_t numColsInRoc = 52;
0443     static constexpr uint16_t lastRowInRoc = numRowsInRoc - 1;
0444     static constexpr uint16_t lastColInRoc = numColsInRoc - 1;
0445 
0446     static constexpr uint16_t numRowsInModule = 2 * numRowsInRoc;
0447     static constexpr uint16_t numColsInModule = 8 * numColsInRoc;
0448     static constexpr uint16_t lastRowInModule = numRowsInModule - 1;
0449     static constexpr uint16_t lastColInModule = numColsInModule - 1;
0450 
0451     static constexpr uint16_t clusterBinning = numColsInModule + 2;
0452     static constexpr uint16_t clusterBits = 9;
0453 
0454     static constexpr uint16_t numberOfModulesInBarrel = 1184;
0455     static constexpr uint16_t numberOfModulesInLadder = 8;
0456     static constexpr uint16_t numberOfLaddersInBarrel = numberOfModulesInBarrel / numberOfModulesInLadder;
0457 
0458     static constexpr uint16_t firstEndcapPos = 4;
0459     static constexpr uint16_t firstEndcapNeg = 7;
0460 
0461     static constexpr int16_t xOffset = -81;
0462 
0463     static constexpr char const *nameModifier = "";
0464 
0465     static constexpr uint32_t const *layerStart = phase1PixelTopology::layerStart;
0466     static constexpr float const *minz = phase1PixelTopology::minz;
0467     static constexpr float const *maxz = phase1PixelTopology::maxz;
0468     static constexpr float const *maxr = phase1PixelTopology::maxr;
0469 
0470     static constexpr uint8_t const *layerPairs = phase1PixelTopology::layerPairs;
0471     static constexpr int16_t const *phicuts = phase1PixelTopology::phicuts;
0472 
0473     static constexpr inline bool isEdgeX(uint16_t px) { return (px == 0) | (px == lastRowInModule); }
0474 
0475     static constexpr inline bool isEdgeY(uint16_t py) { return (py == 0) | (py == lastColInModule); }
0476 
0477     static constexpr inline uint16_t toRocX(uint16_t px) { return (px < numRowsInRoc) ? px : px - numRowsInRoc; }
0478 
0479     static constexpr inline uint16_t toRocY(uint16_t py) {
0480       auto roc = divu52(py);
0481       return py - 52 * roc;
0482     }
0483 
0484     static constexpr inline bool isBigPixX(uint16_t px) { return (px == 79) | (px == 80); }
0485     static constexpr inline bool isBigPixY(uint16_t py) {
0486       auto ly = toRocY(py);
0487       return (ly == 0) | (ly == lastColInRoc);
0488     }
0489 
0490     static constexpr inline uint16_t localX(uint16_t px) {
0491       auto shift = 0;
0492       if (px > lastRowInRoc)
0493         shift += 1;
0494       if (px > numRowsInRoc)
0495         shift += 1;
0496       return px + shift;
0497     }
0498 
0499     static constexpr inline uint16_t localY(uint16_t py) {
0500       auto roc = divu52(py);
0501       auto shift = 2 * roc;
0502       auto yInRoc = py - 52 * roc;
0503       if (yInRoc > 0)
0504         shift += 1;
0505       return py + shift;
0506     }
0507   };
0508 
0509   template <typename T>
0510   using isPhase1Topology = typename std::enable_if<std::is_base_of<Phase1, T>::value>::type;
0511 
0512   template <typename T>
0513   using isPhase2Topology = typename std::enable_if<std::is_base_of<Phase2, T>::value>::type;
0514 
0515   // struct HIonPhase1 : public Phase1 {
0516   //     static constexpr uint32_t maxNumberOfDoublets=3*1024*1024;};
0517 
0518 }  // namespace pixelTopology
0519 
0520 #endif  // Geometry_CommonTopologies_SimplePixelTopology_h