GPU: pass the work-item indices to the helpers that lack them

MSL has no ambient work-item builtins: the thread and threadgroup positions exist
only as attributes on the kernel entry point, so get_local_id() and its siblings
cannot read them from a device function the way CUDA's threadIdx or OpenCL's
get_local_id() can.

Almost every call site is already inside a function that receives nBlocks,
nThreads, iBlock and iThread, which is how the CPU backend has always worked:
four of the six helpers expand to the bare iBlock and nBlocks there, so they only
compile where those are in scope. Four functions have no index in scope at all
and get them passed in: sortInBlock, GPUTPCCFClusterizer::buildCluster,
GPUTPCCFNoiseSuppression::findMinimaAndPeaks and GPUTPCCFPeakFinder::isPeak.

GPUCA_THREAD_INFO_DECL and GPUCA_THREAD_INFO_PROVIDE gate that so only Metal pays
for it: both expand to nothing on every other backend, and the helpers keep
reading get_local_id(0) rather than open-coding the index arithmetic.

Author: ktf

GPU/Common/GPUCommonAlgorithm.h

@@ -32,13 +32,13 @@ class GPUCommonAlgorithm
   template <class T>
   GPUd() static void sort(T* begin, T* end);
   template <class T>
-  GPUd() static void sortInBlock(T* begin, T* end);
+  GPUd() static void sortInBlock(T* begin, T* end GPUCA_THREAD_INFO_DECL);
   template <class T>
   GPUd() static void sortDeviceDynamic(T* begin, T* end);
   template <class T, class S>
   GPUd() static void sort(T* begin, T* end, const S& comp);
   template <class T, class S>
-  GPUd() static void sortInBlock(T* begin, T* end, const S& comp);
+  GPUd() static void sortInBlock(T* begin, T* end, const S& comp GPUCA_THREAD_INFO_DECL);
   template <class T, class S>
   GPUd() static void sortDeviceDynamic(T* begin, T* end, const S& comp);
 #if __cplusplus >= 202002L // sortOnDevice takes an auto parameter
@@ -268,17 +268,17 @@ GPUdi() void GPUCommonAlgorithm::sort(T* begin, T* end, const S& comp)
 }
 
 template <class T>
-GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end)
+GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end GPUCA_THREAD_INFO_DECL)
 {
 #ifndef GPUCA_GPUCODE
   GPUCommonAlgorithm::sort(begin, end);
 #else
-  GPUCommonAlgorithm::sortInBlock(begin, end, [](auto&& x, auto&& y) { return x < y; });
+  GPUCommonAlgorithm::sortInBlock(begin, end, [](auto&& x, auto&& y) { return x < y; } GPUCA_THREAD_INFO_PROVIDE);
 #endif
 }
 
 template <class T, class S>
-GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end, const S& comp)
+GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end, const S& comp GPUCA_THREAD_INFO_DECL)
 {
 #ifndef GPUCA_GPUCODE
   GPUCommonAlgorithm::sort(begin, end, comp);

GPU/Common/GPUCommonAlgorithmThrust.h

@@ -63,15 +63,15 @@ GPUdi() void GPUCommonAlgorithm::sort(T* begin, T* end, const S& comp)
 }
 
 template <class T>
-GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end) // TODO: Try cub::BlockMergeSort
+GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end GPUCA_THREAD_INFO_DECL) // TODO: Try cub::BlockMergeSort
 {
   if (get_local_id(0) == 0) {
     sortDeviceDynamic(begin, end);
   }
 }
 
 template <class T, class S>
-GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end, const S& comp)
+GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end, const S& comp GPUCA_THREAD_INFO_DECL)
 {
   if (get_local_id(0) == 0) {
     sortDeviceDynamic(begin, end, comp);

GPU/Common/GPUCommonDefAPI.h

@@ -286,5 +286,16 @@
   #define get_group_id(dim) iBlock
 #endif
 
+// A device function that uses the helpers above but has none of the indices in
+// scope needs them passed in on Metal, where they are ordinary parameters.
+// Both expand to nothing everywhere else, and go last in the parameter list.
+#ifdef __METAL__
+  #define GPUCA_THREAD_INFO_DECL , int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread
+  #define GPUCA_THREAD_INFO_PROVIDE , nBlocks, nThreads, iBlock, iThread
+#else
+  #define GPUCA_THREAD_INFO_DECL
+  #define GPUCA_THREAD_INFO_PROVIDE
+#endif
+
 // clang-format on
 #endif

GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx

@@ -274,16 +274,16 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step1un
         static_assert(GPUCA_GET_THREAD_COUNT(GPUCA_LB_GPUTPCCompressionKernels_step1unattached) * 2 <= constants::TPC_COMP_CHUNK_SIZE);
 #endif
 #ifdef GPUCA_DETERMINISTIC_MODE
-        CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZPadTime>(clusters->clusters[iSector][iRow]));
+        CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZPadTime>(clusters->clusters[iSector][iRow]) GPUCA_THREAD_INFO_PROVIDE);
 #else  // GPUCA_DETERMINISTIC_MODE
         if (param.rec.tpc.compressionSortOrder == GPUSettings::SortZPadTime) {
-          CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZPadTime>(clusters->clusters[iSector][iRow]));
+          CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZPadTime>(clusters->clusters[iSector][iRow]) GPUCA_THREAD_INFO_PROVIDE);
         } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortZTimePad) {
-          CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZTimePad>(clusters->clusters[iSector][iRow]));
+          CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZTimePad>(clusters->clusters[iSector][iRow]) GPUCA_THREAD_INFO_PROVIDE);
         } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortPad) {
-          CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortPad>(clusters->clusters[iSector][iRow]));
+          CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortPad>(clusters->clusters[iSector][iRow]) GPUCA_THREAD_INFO_PROVIDE);
         } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortTime) {
-          CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortTime>(clusters->clusters[iSector][iRow]));
+          CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortTime>(clusters->clusters[iSector][iRow]) GPUCA_THREAD_INFO_PROVIDE);
         }
 #endif // GPUCA_DETERMINISTIC_MODE
         GPUbarrier();

GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx

@@ -773,7 +773,7 @@ GPUd() void GPUTPCCFHIPTailConnector::Thread<0>(int32_t nBlocks, int32_t nThread
     } else {
       return t1.qMax < t2.qMax;
     }
-  });
+  } GPUCA_THREAD_INFO_PROVIDE);
   if (iThread > 0) {
     return;
   }

GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.h

@@ -59,7 +59,7 @@ class GPUTPCCFClusterizer : public GPUKernelTemplate
 
   static GPUd() void computeClustersImpl(int32_t, int32_t, int32_t, int32_t, processorType&, const CfFragment&, GPUSharedMemory&, const CfArray2D<PackedCharge>&, const CfChargePos*, const GPUSettingsRec&, MCLabelAccumulator*, uint32_t, uint32_t, uint32_t*, tpc::ClusterNative*, uint32_t*, int8_t);
 
-  static GPUd() void buildCluster(const GPUSettingsRec&, const CfArray2D<PackedCharge>&, CfChargePos, CfChargePos*, PackedCharge*, uint8_t*, ClusterAccumulator*, MCLabelAccumulator*);
+  static GPUd() void buildCluster(const GPUSettingsRec&, const CfArray2D<PackedCharge>&, CfChargePos, CfChargePos*, PackedCharge*, uint8_t*, ClusterAccumulator*, MCLabelAccumulator* GPUCA_THREAD_INFO_DECL);
 
   static GPUd() uint32_t sortIntoBuckets(processorType&, const tpc::ClusterNative&, uint32_t, uint32_t, uint32_t*, tpc::ClusterNative*);
 

GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc

@@ -49,7 +49,7 @@ GPUdii() void GPUTPCCFClusterizer::computeClustersImpl(int32_t nBlocks, int32_t
     smem.buf,
     smem.innerAboveThreshold,
     &pc,
-    labelAcc);
+    labelAcc GPUCA_THREAD_INFO_PROVIDE);
 
   if (idx >= clusternum) {
     return;
@@ -151,7 +151,7 @@ GPUdii() void GPUTPCCFClusterizer::buildCluster(
   PackedCharge* buf,
   uint8_t* innerAboveThreshold,
   ClusterAccumulator* myCluster,
-  MCLabelAccumulator* labelAcc)
+  MCLabelAccumulator* labelAcc GPUCA_THREAD_INFO_DECL)
 {
   uint16_t ll = get_local_id(0);
 

GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.cxx

@@ -60,7 +60,7 @@ GPUdii() void GPUTPCCFNoiseSuppression::noiseSuppressionImpl(int32_t nBlocks, in
     smem.buf,
     &minimas,
     &bigger,
-    &peaksAround);
+    &peaksAround GPUCA_THREAD_INFO_PROVIDE);
 
   peaksAround &= bigger;
 
@@ -173,7 +173,7 @@ GPUd() void GPUTPCCFNoiseSuppression::findMinimaAndPeaks(
   PackedCharge* buf,
   uint64_t* minimas,
   uint64_t* bigger,
-  uint64_t* peaks)
+  uint64_t* peaks GPUCA_THREAD_INFO_DECL)
 {
   uint16_t ll = get_local_id(0);
 

GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.h

@@ -69,7 +69,7 @@ class GPUTPCCFNoiseSuppression : public GPUKernelTemplate
 
   static GPUdi() bool keepPeak(uint64_t, uint64_t);
 
-  static GPUd() void findMinimaAndPeaks(const CfArray2D<PackedCharge>&, const CfArray2D<uint8_t>&, const GPUSettingsRec&, float, const CfChargePos&, CfChargePos*, PackedCharge*, uint64_t*, uint64_t*, uint64_t*);
+  static GPUd() void findMinimaAndPeaks(const CfArray2D<PackedCharge>&, const CfArray2D<uint8_t>&, const GPUSettingsRec&, float, const CfChargePos&, CfChargePos*, PackedCharge*, uint64_t*, uint64_t*, uint64_t* GPUCA_THREAD_INFO_DECL);
 };
 
 } // namespace o2::gpu

GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.cxx

@@ -38,7 +38,7 @@ GPUdii() bool GPUTPCCFPeakFinder::isPeak(
   const CfArray2D<PackedCharge>& chargeMap,
   const GPUSettingsRec& calib,
   CfChargePos* posBcast,
-  PackedCharge* buf)
+  PackedCharge* buf GPUCA_THREAD_INFO_DECL)
 {
   uint16_t ll = get_local_id(0);
 
@@ -111,7 +111,7 @@ GPUd() void GPUTPCCFPeakFinder::findPeaksImpl(int32_t nBlocks, int32_t nThreads,
   bool hasLostBaseline = pos.valid() ? padHasLostBaseline[pos.gpad] : true;
   charge = hasLostBaseline ? 0.f : charge;
 
-  uint8_t peak = isPeak(smem, charge, pos, SCRATCH_PAD_SEARCH_N, chargeMap, calib, smem.posBcast, smem.buf);
+  uint8_t peak = isPeak(smem, charge, pos, SCRATCH_PAD_SEARCH_N, chargeMap, calib, smem.posBcast, smem.buf GPUCA_THREAD_INFO_PROVIDE);
 
   // Exit early if dummy. See comment above.
   bool iamDummy = (idx >= digitnum);