From a1ae68e45147000afa2187c6f68b6688909a4783 Mon Sep 17 00:00:00 2001 From: ALICE Action Bot Date: Mon, 21 Sep 2026 09:09:13 +0000 Subject: [PATCH] Please consider the following formatting changes --- GPU/Common/GPUCommonDouble.h | 27 +++++++++---------- GPU/Common/GPUCommonMath.h | 10 +++++-- .../Base/GPUReconstructionKernelMacros.h | 2 +- .../Base/metal/GPUReconstructionMetal.mm | 3 ++- .../SectorTracker/GPUTPCTrackParam.cxx | 2 +- .../TPCClusterFinder/CfChargePos.h | 3 ++- 6 files changed, 27 insertions(+), 20 deletions(-) diff --git a/GPU/Common/GPUCommonDouble.h b/GPU/Common/GPUCommonDouble.h index 1241d97e90c09..7b291df544a64 100644 --- a/GPU/Common/GPUCommonDouble.h +++ b/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(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/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