From d2a8517f5c28030f5d4061f7db27a6404a58aa8d Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Fri, 9 Oct 2026 13:17:18 +0200 Subject: [PATCH] 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. --- GPU/Common/GPUCommonAlgorithm.h | 10 +++++----- GPU/Common/GPUCommonAlgorithmThrust.h | 4 ++-- GPU/Common/GPUCommonDefAPI.h | 11 +++++++++++ .../DataCompression/GPUTPCCompressionKernels.cxx | 10 +++++----- .../TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx | 2 +- .../TPCClusterFinder/GPUTPCCFClusterizer.h | 2 +- .../TPCClusterFinder/GPUTPCCFClusterizer.inc | 4 ++-- .../TPCClusterFinder/GPUTPCCFNoiseSuppression.cxx | 4 ++-- .../TPCClusterFinder/GPUTPCCFNoiseSuppression.h | 2 +- .../TPCClusterFinder/GPUTPCCFPeakFinder.cxx | 4 ++-- GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.h | 2 +- .../TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx | 8 ++++---- 12 files changed, 37 insertions(+), 26 deletions(-) diff --git a/GPU/Common/GPUCommonAlgorithm.h b/GPU/Common/GPUCommonAlgorithm.h index 338a6fb4c5ad4..ef9b54ef3f96d 100644 --- a/GPU/Common/GPUCommonAlgorithm.h +++ b/GPU/Common/GPUCommonAlgorithm.h @@ -32,13 +32,13 @@ class GPUCommonAlgorithm template GPUd() static void sort(T* begin, T* end); template - GPUd() static void sortInBlock(T* begin, T* end); + GPUd() static void sortInBlock(T* begin, T* end GPUCA_THREAD_INFO_DECL); template GPUd() static void sortDeviceDynamic(T* begin, T* end); template GPUd() static void sort(T* begin, T* end, const S& comp); template - 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 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 -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 -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); diff --git a/GPU/Common/GPUCommonAlgorithmThrust.h b/GPU/Common/GPUCommonAlgorithmThrust.h index 7af3138d45490..8c152748bf147 100644 --- a/GPU/Common/GPUCommonAlgorithmThrust.h +++ b/GPU/Common/GPUCommonAlgorithmThrust.h @@ -63,7 +63,7 @@ GPUdi() void GPUCommonAlgorithm::sort(T* begin, T* end, const S& comp) } template -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); @@ -71,7 +71,7 @@ GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end) // TODO: Try cub: } template -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); diff --git a/GPU/Common/GPUCommonDefAPI.h b/GPU/Common/GPUCommonDefAPI.h index 36e8cd403e083..357d304550513 100644 --- a/GPU/Common/GPUCommonDefAPI.h +++ b/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 diff --git a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx index b499ea10e679b..e208f6ba44b9f 100644 --- a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx +++ b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx @@ -274,16 +274,16 @@ GPUdii() void GPUTPCCompressionKernels::Thread(clusters->clusters[iSector][iRow])); + CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(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(clusters->clusters[iSector][iRow])); + CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow]) GPUCA_THREAD_INFO_PROVIDE); } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortZTimePad) { - CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow])); + CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow]) GPUCA_THREAD_INFO_PROVIDE); } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortPad) { - CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow])); + CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow]) GPUCA_THREAD_INFO_PROVIDE); } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortTime) { - CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow])); + CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow]) GPUCA_THREAD_INFO_PROVIDE); } #endif // GPUCA_DETERMINISTIC_MODE GPUbarrier(); diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx index 1696defc1089d..8062d6a878100 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx +++ b/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; } diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.h b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.h index ce673c778e42d..b04b37bc64975 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.h +++ b/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&, 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&, CfChargePos, CfChargePos*, PackedCharge*, uint8_t*, ClusterAccumulator*, MCLabelAccumulator*); + static GPUd() void buildCluster(const GPUSettingsRec&, const CfArray2D&, 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*); diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc index ca396f8aab83e..aeacd606f6197 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc +++ b/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); diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.cxx index 4dfa50d9439e4..1237e1fe5930f 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.cxx +++ b/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); diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.h b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.h index bdee75dc87732..f07f0d7cc2066 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.h +++ b/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&, const CfArray2D&, const GPUSettingsRec&, float, const CfChargePos&, CfChargePos*, PackedCharge*, uint64_t*, uint64_t*, uint64_t*); + static GPUd() void findMinimaAndPeaks(const CfArray2D&, const CfArray2D&, const GPUSettingsRec&, float, const CfChargePos&, CfChargePos*, PackedCharge*, uint64_t*, uint64_t*, uint64_t* GPUCA_THREAD_INFO_DECL); }; } // namespace o2::gpu diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.cxx index 7c93435f8bef8..d0b134badc49a 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.cxx @@ -38,7 +38,7 @@ GPUdii() bool GPUTPCCFPeakFinder::isPeak( const CfArray2D& 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); diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.h b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.h index 0d61378d3e6f2..93727319b86a8 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.h +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.h @@ -53,7 +53,7 @@ class GPUTPCCFPeakFinder : public GPUKernelTemplate private: static GPUd() void findPeaksImpl(int32_t, int32_t, int32_t, int32_t, GPUSharedMemory&, const CfArray2D&, const uint8_t*, const CfChargePos*, tpccf::SizeT, const GPUSettingsRec&, const TPCPadGainCalib&, uint8_t*, CfArray2D&); - static GPUd() bool isPeak(GPUSharedMemory&, tpccf::Charge, const CfChargePos&, uint16_t, const CfArray2D&, const GPUSettingsRec&, CfChargePos*, PackedCharge*); + static GPUd() bool isPeak(GPUSharedMemory&, tpccf::Charge, const CfChargePos&, uint16_t, const CfArray2D&, const GPUSettingsRec&, CfChargePos*, PackedCharge* GPUCA_THREAD_INFO_DECL); }; } // namespace o2::gpu diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx index 902c04e6e92f3..7afc5079e7f4c 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx @@ -344,7 +344,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Threadfragment).isOverlap(peak.time())) { if (clusterer.mPclusterPosInRow) { @@ -538,7 +538,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Threadfragment).isOverlap(peak.time())) { if (clusterer.mPclusterPosInRow) {