Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
davidrohr merged 1 commit into from
Oct 10, 2026
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 5 additions & 5 deletions GPU/Common/GPUCommonAlgorithm.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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);
Expand Down
4 changes: 2 additions & 2 deletions GPU/Common/GPUCommonAlgorithmThrust.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
11 changes: 11 additions & 0 deletions GPU/Common/GPUCommonDefAPI.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
10 changes: 5 additions & 5 deletions GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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();
Expand Down
2 changes: 1 addition & 1 deletion GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
}
Expand Down
2 changes: 1 addition & 1 deletion GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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*);

Expand Down
4 changes: 2 additions & 2 deletions GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down Expand Up @@ -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);

Expand Down
4 changes: 2 additions & 2 deletions GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -60,7 +60,7 @@ GPUdii() void GPUTPCCFNoiseSuppression::noiseSuppressionImpl(int32_t nBlocks, in
smem.buf,
&minimas,
&bigger,
&peaksAround);
&peaksAround GPUCA_THREAD_INFO_PROVIDE);

peaksAround &= bigger;

Expand Down Expand Up @@ -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);

Expand Down
2 changes: 1 addition & 1 deletion GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
4 changes: 2 additions & 2 deletions GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);

Expand Down Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,7 +53,7 @@ class GPUTPCCFPeakFinder : public GPUKernelTemplate
private:
static GPUd() void findPeaksImpl(int32_t, int32_t, int32_t, int32_t, GPUSharedMemory&, const CfArray2D<PackedCharge>&, const uint8_t*, const CfChargePos*, tpccf::SizeT, const GPUSettingsRec&, const TPCPadGainCalib&, uint8_t*, CfArray2D<uint8_t>&);

static GPUd() bool isPeak(GPUSharedMemory&, tpccf::Charge, const CfChargePos&, uint16_t, const CfArray2D<PackedCharge>&, const GPUSettingsRec&, CfChargePos*, PackedCharge*);
static GPUd() bool isPeak(GPUSharedMemory&, tpccf::Charge, const CfChargePos&, uint16_t, const CfArray2D<PackedCharge>&, const GPUSettingsRec&, CfChargePos*, PackedCharge* GPUCA_THREAD_INFO_DECL);
};

} // namespace o2::gpu
Expand Down
8 changes: 4 additions & 4 deletions GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -344,7 +344,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread<GPUTPCNNClusterizerKernels::pub
smem.buf,
smem.innerAboveThreshold,
&dummy_pc,
labelAcc);
labelAcc GPUCA_THREAD_INFO_PROVIDE);
}
return;
}
Expand All @@ -364,7 +364,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread<GPUTPCNNClusterizerKernels::pub
smem.buf,
smem.innerAboveThreshold,
&dummy_pc,
labelAcc);
labelAcc GPUCA_THREAD_INFO_PROVIDE);
}
if ((clusterer.mPmemory->fragment).isOverlap(peak.time())) {
if (clusterer.mPclusterPosInRow) {
Expand Down Expand Up @@ -538,7 +538,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread<GPUTPCNNClusterizerKernels::pub
smem.buf,
smem.innerAboveThreshold,
&dummy_pc,
labelAcc);
labelAcc GPUCA_THREAD_INFO_PROVIDE);
}
return;
}
Expand All @@ -558,7 +558,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread<GPUTPCNNClusterizerKernels::pub
smem.buf,
smem.innerAboveThreshold,
&dummy_pc,
labelAcc);
labelAcc GPUCA_THREAD_INFO_PROVIDE);
}
if ((clusterer.mPmemory->fragment).isOverlap(peak.time())) {
if (clusterer.mPclusterPosInRow) {
Expand Down
Loading