ktf · alibuild · Sep 21, 2026
diff --git 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<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

diff --git a/GPU/Common/GPUCommonMath.h b/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>

diff --git 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
@@ -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
@@ -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
@@ -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