GPU: Apply clang-format or suppress it

Author: davidrohr

.clang-format

@@ -27,6 +27,7 @@ IndentWidth:     2
 # It is broken on windows. Breaks all #include "header.h"
 ---
 Language:        Cpp
+AttributeMacros: ['GPUrestrict'] # GPU qualifier macro: GPUrestrict() is otherwise parsed as function call, e.g. 'a -> b' after it
 MaxEmptyLinesToKeep: 1
 KeepEmptyLinesAtTheStartOfBlocks: true
 NamespaceIndentation: None

GPU/Common/GPUCommonArray.h

@@ -41,8 +41,10 @@ struct array {
   }
   T m_internal_V__[N];
 };
+// clang-format off
 template <class T, class... E>
 GPUd() array(T, E...)->array<T, 1 + sizeof...(E)>;
+// clang-format on
 #else
 template <typename T, size_t N>
 using array = std::array<T, N>;

GPU/Common/GPUCommonConstants.h

@@ -21,6 +21,6 @@ namespace o2::gpu::gpu_common_constants
 {
 static GPUglobalconstexpr() const float kCLight = 0.000299792458f; // TODO: Duplicate of MathConstants, fix this now that we use only OpenCL CPP
 static GPUglobalconstexpr() const float kZeroFieldCut = 0.013f;
-}
+} // namespace o2::gpu::gpu_common_constants
 
 #endif

GPU/GPUTracking/Base/GPUReconstructionKernelIncludes.template.h

@@ -14,5 +14,4 @@
 
 // clang-format off
 $<JOIN:$<LIST:TRANSFORM,$<LIST:TRANSFORM,$<LIST:REMOVE_DUPLICATES,$<TARGET_PROPERTY:O2_GPU_KERNELS,O2_GPU_KERNEL_INCLUDES>>,APPEND,.h">,PREPEND,#include ">,
->
-// clang-format on
+> // clang-format on

GPU/GPUTracking/Base/GPUReconstructionKernelList.template.h

@@ -16,5 +16,4 @@
 #include "GPUReconstructionKernelMacros.h"
 
 // clang-format off
-$<JOIN:$<TARGET_PROPERTY:O2_GPU_KERNELS,O2_GPU_KERNELS>,>
-// clang-format on
+$<JOIN:$<TARGET_PROPERTY:O2_GPU_KERNELS,O2_GPU_KERNELS>,> // clang-format on

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu

@@ -331,9 +331,9 @@ int32_t GPUReconstructionCUDA::InitDevice_Runtime()
       }
     }
 
-#ifndef __HIPCC__ // CUDA
+#ifndef __HIPCC__                                                      // CUDA
     dummyInitKernel<<<mMultiprocessorCount, 256>>>(mDeviceMemoryBase); // TODO: Can't we just use the CUDA version and hipify will take care of the rest?
-#else // HIP
+#else                                                                  // HIP
     hipLaunchKernelGGL(HIP_KERNEL_NAME(dummyInitKernel), dim3(mMultiprocessorCount), dim3(256), 0, 0, mDeviceMemoryBase);
 #endif
 

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAkernel.template.cu

@@ -21,10 +21,7 @@
 
 // clang-format off
 @O2_GPU_KERNEL_TEMPLATE_FILES@
-// clang-format on
 
 extern "C" {
-// clang-format off
 @O2_GPU_KERNEL_TEMPLATE_REPLACE@
-// clang-format on
-}
+} // clang-format on: not before the }, the placeholders confuse clang-format

GPU/GPUTracking/Base/hip/GPUReconstructionHIPkernel.template.hip

@@ -21,10 +21,7 @@
 
 // clang-format off
 @O2_GPU_KERNEL_TEMPLATE_FILES@
-// clang-format on
 
 extern "C" {
-// clang-format off
 @O2_GPU_KERNEL_TEMPLATE_REPLACE@
-// clang-format on
-}
+} // clang-format on: not before the }, the placeholders confuse clang-format

GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx

@@ -67,7 +67,7 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step0at
       if (!(param.rec.tpc.compressionTypeMask & GPUSettings::CompressionTrackModel)) {
         continue; // No track model compression
       }
-      const ClusterNative& GPUrestrict() orgCl = clusters -> clusters[hit.sector][hit.row][hit.num - clusters->clusterOffset[hit.sector][hit.row]];
+      const ClusterNative& GPUrestrict() orgCl = clusters->clusters[hit.sector][hit.row][hit.num - clusters->clusterOffset[hit.sector][hit.row]];
       constexpr GPUTPCGeometry geo;
       float x = geo.Row2X(hit.row);
       float y = track.LinearPad2Y(hit.sector, orgCl.getPad(), geo.PadWidth(hit.row), geo.NPads(hit.row));
@@ -291,7 +291,7 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step1un
 
       for (uint32_t j = get_local_id(0); j < count; j += get_local_size(0)) {
         int32_t outidx = idOffsetOut + totalCount + j;
-        const ClusterNative& GPUrestrict() orgCl = clusters -> clusters[iSector][iRow][sortBuffer[j]];
+        const ClusterNative& GPUrestrict() orgCl = clusters->clusters[iSector][iRow][sortBuffer[j]];
 
         int32_t preId = j != 0 ? (int32_t)sortBuffer[j - 1] : (totalCount != 0 ? (int32_t)smem.lastIndex : -1);
         GPUTPCCompression_EncodeUnattached(param.rec.tpc.compressionTypeMask, orgCl, c.timeDiffU[outidx], c.padDiffU[outidx], preId == -1 ? nullptr : &clusters->clusters[iSector][iRow][preId]);

GPU/GPUTracking/DataTypes/GPUTPCGeometry.h

@@ -121,7 +121,7 @@ class GPUTPCGeometry
   GPUd() static constexpr int32_t MaxNPadsPerRow() { return 138; }
   GPUd() static constexpr float PadWidth(int32_t row) { return (gputpcgeometry_internal::mPadWidthRow[row]); }
 #else
-  GPUd() static constexpr int32_t GetRegion(int32_t row) { return (row < 63 ? 0 : row < 63 + 64 ? 1 : 2); }
+  GPUd() static constexpr int32_t GetRegion(int32_t row) { return (row < 63 ? 0 : (row < 63 + 64 ? 1 : 2)); }
   GPUd() static constexpr int32_t GetRegionRows(int32_t region) { return 0; }  // dummy
   GPUd() static constexpr int32_t GetRegionStart(int32_t region) { return 0; } // dummy
   GPUd() static constexpr int32_t GetROC(int32_t row) { return GetRegion(row); }