Please consider the following formatting changes to #15817

Common/Constants/include/CommonConstants/LHCConstants.h

@@ -32,7 +32,7 @@ enum BeamDirection : int { BeamA, // beamA = beam 0,
                            NBeamDirections,
                            InteractingBC = -1 // as used in the BunchFilling class
 };
-GPUglobalconstexpr() int LHCMaxBunches = 3564;                              // max N bunches
+GPUglobalconstexpr() int LHCMaxBunches = 3564;                                               // max N bunches
 GPUglobalconstexpr() o2::gpu::GPUdoubleValue LHCRFFreq = 400.789e6;                          // LHC RF frequency in Hz
 GPUglobalconstexpr() o2::gpu::GPUdoubleValue LHCBunchSpacingNS = 10 * 1.e9 / LHCRFFreq;      // bunch spacing in ns (10 RFbuckets)
 GPUglobalconstexpr() o2::gpu::GPUdoubleValue LHCOrbitNS = LHCMaxBunches * LHCBunchSpacingNS; // orbit duration in ns

DataFormats/Detectors/TRD/include/DataFormatsTRD/Constants.h

@@ -40,27 +40,27 @@ GPUglobalconstexpr() int NFLP = 12;             ///< the number of FLP we have.
 GPUglobalconstexpr() int NCRUPERFLP = 3;        ///< the number of CRU per FLP
 GPUglobalconstexpr() int TRDLINKID = 15;        ///< hard coded link id, specific to TRD
 
-GPUglobalconstexpr() int NCOLUMN = 144; ///< the number of pad columns for each chamber
-GPUglobalconstexpr() int NROWC0 = 12;   ///< the number of pad rows for chambers of type C0 (installed in stack 2)
-GPUglobalconstexpr() int NROWC1 = 16;   ///< the number of pad rows for chambers of type C1 (installed in stacks 0, 1, 3 and 4)
+GPUglobalconstexpr() int NCOLUMN = 144;                          ///< the number of pad columns for each chamber
+GPUglobalconstexpr() int NROWC0 = 12;                            ///< the number of pad rows for chambers of type C0 (installed in stack 2)
+GPUglobalconstexpr() int NROWC1 = 16;                            ///< the number of pad rows for chambers of type C1 (installed in stacks 0, 1, 3 and 4)
 GPUglobalconstexpr() int FIRSTROW[NSTACK] = {0, 16, 32, 44, 60}; ///< first pad row for each stack
 
-GPUglobalconstexpr() int NMCMROB = 16;     ///< the number of MCMs per ROB
-GPUglobalconstexpr() int NMCMHCMAX = 64;   ///< the maximum number of MCMs for one half chamber (C1 type)
-GPUglobalconstexpr() int NMCMROBINROW = 4; ///< the number of MCMs per ROB in row direction
-GPUglobalconstexpr() int NMCMROBINCOL = 4; ///< the number of MCMs per ROB in column direction
-GPUglobalconstexpr() int NROBC0 = 6;       ///< the number of ROBs per C0 chamber
-GPUglobalconstexpr() int NROBC1 = 8;       ///< the number of ROBs per C1 chamber
-GPUglobalconstexpr() int NADCMCM = 21;     ///< the number of ADC channels per MCM
-GPUglobalconstexpr() int NCOLMCM = 18;     ///< the number of pads per MCM
+GPUglobalconstexpr() int NMCMROB = 16;                                                                                    ///< the number of MCMs per ROB
+GPUglobalconstexpr() int NMCMHCMAX = 64;                                                                                  ///< the maximum number of MCMs for one half chamber (C1 type)
+GPUglobalconstexpr() int NMCMROBINROW = 4;                                                                                ///< the number of MCMs per ROB in row direction
+GPUglobalconstexpr() int NMCMROBINCOL = 4;                                                                                ///< the number of MCMs per ROB in column direction
+GPUglobalconstexpr() int NROBC0 = 6;                                                                                      ///< the number of ROBs per C0 chamber
+GPUglobalconstexpr() int NROBC1 = 8;                                                                                      ///< the number of ROBs per C1 chamber
+GPUglobalconstexpr() int NADCMCM = 21;                                                                                    ///< the number of ADC channels per MCM
+GPUglobalconstexpr() int NCOLMCM = 18;                                                                                    ///< the number of pads per MCM
 GPUglobalconstexpr() int NCHANNELSPERROW = NMCMROBINCOL * 2 * NADCMCM;                                                    ///< the number of readout channels per pad row
 GPUglobalconstexpr() int NCHANNELSC0 = NROWC0 * NCHANNELSPERROW;                                                          ///< the number of readout channels per C0 chamber
 GPUglobalconstexpr() int NCHANNELSC1 = NROWC1 * NCHANNELSPERROW;                                                          ///< the number of readout channels per C1 chamber
 GPUglobalconstexpr() int NCHANNELSTOTAL = NSECTOR * NLAYER * (NSTACK - 1) * NCHANNELSC1 + NSECTOR * NLAYER * NCHANNELSC0; ///< the total number of readout channels for TRD
 GPUglobalconstexpr() int NCHANNELSPERSECTOR = NCHANNELSTOTAL / NSECTOR;                                                   ///< then number of readout channels per sector
 GPUglobalconstexpr() int NCHANNELSPERLAYER = NCHANNELSPERSECTOR / NLAYER;                                                 ///< then number of readout channels per layer
-GPUglobalconstexpr() int NCPU = 4;         ///< the number of CPUs inside the TRAP chip
-GPUglobalconstexpr() int NCHARGES = 3;     ///< the number of charges per tracklet (Q0/1/2)
+GPUglobalconstexpr() int NCPU = 4;                                                                                        ///< the number of CPUs inside the TRAP chip
+GPUglobalconstexpr() int NCHARGES = 3;                                                                                    ///< the number of charges per tracklet (Q0/1/2)
 
 // the values below should come out of the TRAP config in the future
 GPUglobalconstexpr() int NBITSTRKLPOS = 11;                                      ///< number of bits for position in tracklet64 word
@@ -73,9 +73,9 @@ GPUglobalconstexpr() float GRANULARITYTRKLSLOPE = 1.f / PADGRANULARITYTRKLSLOPE;
 GPUglobalconstexpr() int ADCBASELINE = 10;                                       ///< baseline in ADC units
 
 // OS: Should this not be flexible for example in case of Kr calib?
-GPUglobalconstexpr() int TIMEBINS = 30;            ///< the number of time bins
-GPUglobalconstexpr() float MAXIMPACTANGLE = 25.f;  ///< the maximum impact angle for tracks relative to the TRD detector plane to be considered for vDrift and ExB calibration
-GPUglobalconstexpr() int NBINSANGLEDIFF = 25;      ///< the number of bins for the track angle used for the vDrift and ExB calibration based on the tracking
+GPUglobalconstexpr() int TIMEBINS = 30;           ///< the number of time bins
+GPUglobalconstexpr() float MAXIMPACTANGLE = 25.f; ///< the maximum impact angle for tracks relative to the TRD detector plane to be considered for vDrift and ExB calibration
+GPUglobalconstexpr() int NBINSANGLEDIFF = 25;     ///< the number of bins for the track angle used for the vDrift and ExB calibration based on the tracking
 #ifndef GPUCA_GPUCODE_DEVICE
 // calibration defaults, host only: these are double and never used in device code
 constexpr double VDRIFTDEFAULT = 1.546; ///< default value for vDrift
@@ -85,9 +85,9 @@ constexpr double EXBDEFAULT = 0.0;      ///< default value for LorentzAngle
 constexpr double EXBMIN = -0.4;         ///< min value for LorentzAngle
 constexpr double EXBMAX = 0.4;          ///< max value for LorentzAngle
 #endif
-GPUglobalconstexpr() int NBINSGAINCALIB = 320;     ///< number of bins in the charge (Q0+Q1+Q2) histogram for gain calibration
-GPUglobalconstexpr() float MPVDEDXDEFAULT = 42.;   ///< default Most Probable Value of TRD dEdx
-GPUglobalconstexpr() float T0DEFAULT = 1.2;        ///< default value for t0
+GPUglobalconstexpr() int NBINSGAINCALIB = 320;   ///< number of bins in the charge (Q0+Q1+Q2) histogram for gain calibration
+GPUglobalconstexpr() float MPVDEDXDEFAULT = 42.; ///< default Most Probable Value of TRD dEdx
+GPUglobalconstexpr() float T0DEFAULT = 1.2;      ///< default value for t0
 
 // array size to store incoming half cru payload.
 GPUglobalconstexpr() int HBFBUFFERMAX = 1048576;                 ///< max buffer size for data read from a half cru, (all events)

DataFormats/Reconstruction/include/ReconstructionDataFormats/PID.h

@@ -102,8 +102,8 @@ class PID
   static GPUglobalconstexpr() ID Alpha = 8;
 
   static GPUglobalconstexpr() ID First = Electron;
-  static GPUglobalconstexpr() ID Last = Alpha;     ///< if extra IDs added, update this !!!
-  static GPUglobalconstexpr() ID NIDs = Last + 1;  ///< number of defined IDs
+  static GPUglobalconstexpr() ID Last = Alpha;    ///< if extra IDs added, update this !!!
+  static GPUglobalconstexpr() ID NIDs = Last + 1; ///< number of defined IDs
 
   // PID for derived particles
   static GPUglobalconstexpr() ID PI0 = 9;

DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackParametrization.h

@@ -100,12 +100,12 @@ enum DirType : int { DirInward = -1,
 GPUglobalconstexpr() int kNParams = 5, kCovMatSize = 15, kLabCovMatSize = 21;
 
 GPUglobalconstexpr() float kCY2max = 100 * 100, // SigmaY<=100cm
-  kCZ2max = 100 * 100,               // SigmaZ<=100cm
-  kCSnp2max = 1 * 1,                 // SigmaSin<=1
-  kCTgl2max = 1 * 1,                 // SigmaTan<=1
-  kC1Pt2max = 100 * 100,             // Sigma1/Pt<=100 1/GeV
-  kMostProbablePt = 0.6f,            // Most Probable Pt (GeV), for running with Bz=0
-  kCalcdEdxAuto = -999.f;            // value indicating request for dedx calculation
+  kCZ2max = 100 * 100,                          // SigmaZ<=100cm
+  kCSnp2max = 1 * 1,                            // SigmaSin<=1
+  kCTgl2max = 1 * 1,                            // SigmaTan<=1
+  kC1Pt2max = 100 * 100,                        // Sigma1/Pt<=100 1/GeV
+  kMostProbablePt = 0.6f,                       // Most Probable Pt (GeV), for running with Bz=0
+  kCalcdEdxAuto = -999.f;                       // value indicating request for dedx calculation
 
 // access to covariance matrix by row and column
 GPUconstexpr() int CovarMap[kNParams][kNParams] = {{0, 1, 3, 6, 10},

Detectors/TRD/base/include/TRDBase/GeometryBase.h

@@ -140,7 +140,7 @@ class GeometryBase
   static GPUglobalconstexpr() float CWSH = 0.3;    ///< Height of additional wacosit ledge on lower frame
 
   static GPUglobalconstexpr() float CPADW = 0.0; ///>Difference of outer chamber width and pad plane width
-  static GPUglobalconstexpr() float RPADW = 1.0; ///<Difference of outer chamber width and pad plane width
+  static GPUglobalconstexpr() float RPADW = 1.0; ///< Difference of outer chamber width and pad plane width
 
   //
   // Thickness of the the material layers
@@ -188,11 +188,11 @@ class GeometryBase
   static GPUglobalconstexpr() float TIME0BASE = 300.65; ///< Base value for calculation of Time-position of pad 0
   // Time-position of pad 0
   static GPUglobalconstexpr() float TIME0[6] = {TIME0BASE + 0 * (CH + VSPACE),
-                                     TIME0BASE + 1 * (CH + VSPACE),
-                                     TIME0BASE + 2 * (CH + VSPACE),
-                                     TIME0BASE + 3 * (CH + VSPACE),
-                                     TIME0BASE + 4 * (CH + VSPACE),
-                                     TIME0BASE + 5 * (CH + VSPACE)};
+                                                TIME0BASE + 1 * (CH + VSPACE),
+                                                TIME0BASE + 2 * (CH + VSPACE),
+                                                TIME0BASE + 3 * (CH + VSPACE),
+                                                TIME0BASE + 4 * (CH + VSPACE),
+                                                TIME0BASE + 5 * (CH + VSPACE)};
 
   static GPUglobalconstexpr() float XTRDBEG = 288.43; ///< X-coordinate in tracking system of begin of TRD mother volume
   static GPUglobalconstexpr() float XTRDEND = 366.33; ///< X-coordinate in tracking system of end of TRD mother volume

GPU/Common/GPUCommonDouble.h

@@ -89,7 +89,6 @@ GPUhdi() GPUdoubleValue GPUdoubleGet(GPUdoubleStore d) { return d; }
 static_assert(sizeof(GPUdoubleStore) == 8, "GPUdoubleStore must match the size of a double");
 static_assert(alignof(GPUdoubleStore) == 8, "GPUdoubleStore must match the alignment of a double");
 
-
 // Compensated two-float arithmetic, value = mHi + mLo, for the intermediates that
 // are deliberately computed in double even when the track itself is float -- the
 // Jacobian and covariance terms in TrackParametrizationWithError::propagateTo and
@@ -221,21 +220,21 @@ GPUhdi() GPUdoubleBinary64 GPUCommonMath::Abs<GPUdoubleBinary64>(GPUdoubleBinary
 #define GPUCA_DOUBLECALC_BINARY64 4
 
 #ifndef GPUCA_DOUBLECALC
-  #if defined(__METAL__) && defined(__FAST_MATH__)
-    // Fast math reassociates the compensation terms away, so the two-float type
-    // would cost 1.5x for the accuracy of a plain float.
-    #define GPUCA_DOUBLECALC GPUCA_DOUBLECALC_FLOAT
-  #elif defined(__METAL__)
-    #define GPUCA_DOUBLECALC GPUCA_DOUBLECALC_TWOFLOAT
-  #else
-    #define GPUCA_DOUBLECALC GPUCA_DOUBLECALC_DOUBLE
-  #endif
+#if defined(__METAL__) && defined(__FAST_MATH__)
+// Fast math reassociates the compensation terms away, so the two-float type
+// would cost 1.5x for the accuracy of a plain float.
+#define GPUCA_DOUBLECALC GPUCA_DOUBLECALC_FLOAT
+#elif defined(__METAL__)
+#define GPUCA_DOUBLECALC GPUCA_DOUBLECALC_TWOFLOAT
+#else
+#define GPUCA_DOUBLECALC GPUCA_DOUBLECALC_DOUBLE
+#endif
 #endif
 
 #if GPUCA_DOUBLECALC == GPUCA_DOUBLECALC_DOUBLE
-  #ifdef __METAL__
-    #error "MSL has no double; GPUCA_DOUBLECALC_DOUBLE cannot be selected for Metal"
-  #endif
+#ifdef __METAL__
+#error "MSL has no double; GPUCA_DOUBLECALC_DOUBLE cannot be selected for Metal"
+#endif
 typedef double GPUdoubleCalc;
 #elif GPUCA_DOUBLECALC == GPUCA_DOUBLECALC_FLOAT
 typedef float GPUdoubleCalc;
@@ -244,7 +243,7 @@ typedef GPUdoubleCalcImpl GPUdoubleCalc;
 #elif GPUCA_DOUBLECALC == GPUCA_DOUBLECALC_BINARY64
 typedef GPUdoubleBinary64 GPUdoubleCalc;
 #else
-  #error "Invalid setting for GPUCA_DOUBLECALC"
+#error "Invalid setting for GPUCA_DOUBLECALC"
 #endif
 
 } // namespace o2::gpu

GPU/Common/GPUCommonMath.h

@@ -483,9 +483,15 @@ GPUhdi() constexpr int32_t GPUCommonMath::Abs<int32_t>(int32_t x)
 // has the same size and alignment; the overloads keep the address space, which a
 // generic pointer would not carry into atomic_*_explicit.
 template <class T>
-GPUdi() threadgroup metal::atomic<T>* GPUCommonMathMetalAtomic(threadgroup T* p) { return reinterpret_cast<threadgroup metal::atomic<T>*>(p); }
+GPUdi() threadgroup metal::atomic<T>* GPUCommonMathMetalAtomic(threadgroup T* p)
+{
+  return reinterpret_cast<threadgroup metal::atomic<T>*>(p);
+}
 template <class T>
-GPUdi() device metal::atomic<T>* GPUCommonMathMetalAtomic(T* p) { return (device metal::atomic<T>*)p; }
+GPUdi() device metal::atomic<T>* GPUCommonMathMetalAtomic(T* p)
+{
+  return (device metal::atomic<T>*)p;
+}
 #endif
 
 template <class S, class T>

GPU/GPUTracking/Base/GPUReconstructionKernelMacros.h

@@ -91,4 +91,4 @@
 #define GPUCA_KRNL_LB(x_class, x_attributes, ...) GPUCA_KRNL(x_class, (REG, (GPUCA_M_CAT(GPUCA_LB_, GPUCA_M_KRNL_NAME(x_class))), GPUCA_M_STRIP(x_attributes)), __VA_ARGS__)
 
 #endif // O2_GPU_GPURECONSTRUCTIONKERNELMACROS_H
-// clang-format on
+  // clang-format on

GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm

@@ -399,7 +399,8 @@
 
   // the section is part of the mapped image, so it outlives the dispatch_data_t
   // and does not have to be copied
-  dispatch_data_t blob = dispatch_data_create(p, sz, nullptr, ^{});
+  dispatch_data_t blob = dispatch_data_create(p, sz, nullptr, ^{
+                                                     });
 
   NSError* error = nil;
   mInternals->library = [mInternals->device newLibraryWithData:blob error:&error];

GPU/GPUTracking/DataTypes/CalibdEdxTrackTopologyPol.h

@@ -197,9 +197,9 @@ class CalibdEdxTrackTopologyPol : public o2::gpu::FlatObject
   /// ================================================================================================
 
  private:
-  GPUglobalconstexpr() static int32_t FFits{10};                                          ///< total number of fits: 10 regions * 2 charge types
-  GPUglobalconstexpr() static int32_t FDim{5};                                            ///< dimensions of polynomials
-  GPUglobalconstexpr() static int32_t FDegree{3};                                         ///< degree of polynomials
+  GPUglobalconstexpr() static int32_t FFits { 10 };                            ///< total number of fits: 10 regions * 2 charge types
+  GPUglobalconstexpr() static int32_t FDim { 5 };                              ///< dimensions of polynomials
+  GPUglobalconstexpr() static int32_t FDegree { 3 };                           ///< degree of polynomials
   o2::gpu::NDPiecewisePolynomials<FDim, FDegree, false> mCalibPolsqTot[FFits]; ///< polynomial objects storage for the polynomials for qTot
   o2::gpu::NDPiecewisePolynomials<FDim, FDegree, false> mCalibPolsqMax[FFits]; ///< polynomial objects storage for the polynomials for qMax
   float mScalingFactorsqTot[FFits]{1, 1, 1, 1, 1, 1, 1, 1, 1, 1};              ///< value which is used to scale the result of the polynomial for qTot (can be used for normalization)

GPU/GPUTracking/DataTypes/CalibdEdxTrackTopologySpline.h

@@ -227,8 +227,8 @@ class CalibdEdxTrackTopologySpline : public o2::gpu::FlatObject
 #endif
 
  private:
-  GPUglobalconstexpr() static uint32_t FSplines = 10;                           ///< number of splines stored for each type
-  GPUglobalconstexpr() static int32_t FDimX = 3;                                ///< dimensionality of the splines
+  GPUglobalconstexpr() static uint32_t FSplines = 10;                ///< number of splines stored for each type
+  GPUglobalconstexpr() static int32_t FDimX = 3;                     ///< dimensionality of the splines
   SplineType mCalibSplinesqMax[FSplines];                            ///< spline objects storage for the splines for qMax
   SplineType mCalibSplinesqTot[FSplines];                            ///< spline objects storage for the splines for qTot
   float mMaxTanTheta{2.f};                                           ///< max tanTheta for which the correction is stored

GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.cxx

@@ -304,7 +304,7 @@ GPUd() bool GPUTPCTrackParam::TransportToXWithMaterial(float x, GPUTPCTrackLinea
 {
   //* Transport the track parameters to X=x  taking into account material budget
 
-#ifdef __METAL__ // MSL rejects variables declared static at function scope
+#ifdef __METAL__                      // MSL rejects variables declared static at function scope
   constexpr float kRho = 1.025e-3f;   // [g/cm^3]
   constexpr float kRadLen = 28811.7f; //[cm]
   constexpr float kRadLenInv = 1.f / kRadLen;

GPU/GPUTracking/TPCClusterFinder/CfChargePos.h

@@ -38,7 +38,8 @@ struct CfChargePos {
   // INVALID_CHARGE_POS below lives in the constant address space, which a
   // generic `this` does not reach in MSL.
   constexpr GPUhdi() CfChargePos(tpccf::Row row, tpccf::Pad pad, tpccf::TPCFragmentTime t) constant
-    : gpad(tpcGlobalPadIdx(row, pad)), timePadded(t + GPUCF_PADDING_TIME)
+    : gpad(tpcGlobalPadIdx(row, pad)),
+      timePadded(t + GPUCF_PADDING_TIME)
   {
   }
 #endif