diff --git a/Common/Constants/include/CommonConstants/LHCConstants.h b/Common/Constants/include/CommonConstants/LHCConstants.h index 213c56a2378b8..f095260775d38 100644 --- a/Common/Constants/include/CommonConstants/LHCConstants.h +++ b/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 diff --git a/DataFormats/Detectors/TRD/include/DataFormatsTRD/Constants.h b/DataFormats/Detectors/TRD/include/DataFormatsTRD/Constants.h index 3ef5bdbec9ac0..a304bef503973 100644 --- a/DataFormats/Detectors/TRD/include/DataFormatsTRD/Constants.h +++ b/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) diff --git a/DataFormats/Reconstruction/include/ReconstructionDataFormats/PID.h b/DataFormats/Reconstruction/include/ReconstructionDataFormats/PID.h index fb51a7aff6955..f7854bae1ad45 100644 --- a/DataFormats/Reconstruction/include/ReconstructionDataFormats/PID.h +++ b/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; diff --git a/DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackParametrization.h b/DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackParametrization.h index 87a8a990e6fdb..ee9b3c10e05b7 100644 --- a/DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackParametrization.h +++ b/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}, diff --git a/Detectors/TRD/base/include/TRDBase/GeometryBase.h b/Detectors/TRD/base/include/TRDBase/GeometryBase.h index 251e5f8de5e40..bb19472c2fbde 100644 --- a/Detectors/TRD/base/include/TRDBase/GeometryBase.h +++ b/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; ///(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 diff --git a/GPU/Common/GPUCommonMath.h b/GPU/Common/GPUCommonMath.h index 5d8dd6e99bbe8..ae4263c1e3840 100644 --- a/GPU/Common/GPUCommonMath.h +++ b/GPU/Common/GPUCommonMath.h @@ -483,9 +483,15 @@ GPUhdi() constexpr int32_t GPUCommonMath::Abs(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 -GPUdi() threadgroup metal::atomic* GPUCommonMathMetalAtomic(threadgroup T* p) { return reinterpret_cast*>(p); } +GPUdi() threadgroup metal::atomic* GPUCommonMathMetalAtomic(threadgroup T* p) +{ + return reinterpret_cast*>(p); +} template -GPUdi() device metal::atomic* GPUCommonMathMetalAtomic(T* p) { return (device metal::atomic*)p; } +GPUdi() device metal::atomic* GPUCommonMathMetalAtomic(T* p) +{ + return (device metal::atomic*)p; +} #endif template diff --git a/GPU/GPUTracking/Base/GPUReconstructionKernelMacros.h b/GPU/GPUTracking/Base/GPUReconstructionKernelMacros.h index 0887cadd7338d..34b3bea076ac5 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionKernelMacros.h +++ b/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 diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm index 8448a9d86184b..0fb42452e6a05 100644 --- a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm +++ b/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]; diff --git a/GPU/GPUTracking/DataTypes/CalibdEdxTrackTopologyPol.h b/GPU/GPUTracking/DataTypes/CalibdEdxTrackTopologyPol.h index a1a09a18165ef..8110e32c72391 100644 --- a/GPU/GPUTracking/DataTypes/CalibdEdxTrackTopologyPol.h +++ b/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 mCalibPolsqTot[FFits]; ///< polynomial objects storage for the polynomials for qTot o2::gpu::NDPiecewisePolynomials 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) diff --git a/GPU/GPUTracking/DataTypes/CalibdEdxTrackTopologySpline.h b/GPU/GPUTracking/DataTypes/CalibdEdxTrackTopologySpline.h index f1bcdb4413fa8..bf371d8670357 100644 --- a/GPU/GPUTracking/DataTypes/CalibdEdxTrackTopologySpline.h +++ b/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 diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.cxx index 28a1cff148df6..9671b9d1590d3 100644 --- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackParam.cxx +++ b/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; diff --git a/GPU/GPUTracking/TPCClusterFinder/CfChargePos.h b/GPU/GPUTracking/TPCClusterFinder/CfChargePos.h index 64aa0f6fcfe6d..eb41dac2328ea 100644 --- a/GPU/GPUTracking/TPCClusterFinder/CfChargePos.h +++ b/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