Skip to content

Commit 36b209a

Browse files
ktfktf
committed
GPU: take the work-item indices from the kernel parameters
Metal has no ambient work-item builtins: the thread and threadgroup positions exist only as attributes on the kernel entry point, so get_global_id() and its siblings cannot be expressed in a device function the way CUDA's threadIdx or OpenCL's get_global_id() can. Every one of these call sites is inside a function that already receives nBlocks, nThreads, iBlock and iThread, or is one call away from one, so they now use those directly. The substitution is exact on every backend: CUDA, HIP and OpenCL pass precisely these values into Thread(), and the CPU backend passes nThreads = 1 and iThread = 0, which is what the CPU definitions of the macros already assumed. Four helpers had no index in scope and gain one parameter: sortInBlock, GPUTPCCFClusterizer::buildCluster, GPUTPCCFNoiseSuppression::findMinimaAndPeaks and GPUTPCCFPeakFinder::isPeak. GPUCA_SHARED_CACHE and GPUCA_TBB_KERNEL_LOOP likewise take the indices as arguments rather than capturing them from the expansion context. The get_*() macros remain for the kernel entry point in GPUReconstructionKernelMacros.h, which is the one place where Metal does provide them.
1 parent 992cd42 commit 36b209a

29 files changed

Lines changed: 110 additions & 107 deletions

‎GPU/Common/GPUCommonAlgorithm.h‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -32,13 +32,13 @@ class GPUCommonAlgorithm
3232
template <class T>
3333
GPUd() static void sort(T* begin, T* end);
3434
template <class T>
35-
GPUd() static void sortInBlock(T* begin, T* end);
35+
GPUd() static void sortInBlock(int32_t nThreads, int32_t iThread, T* begin, T* end);
3636
template <class T>
3737
GPUd() static void sortDeviceDynamic(T* begin, T* end);
3838
template <class T, class S>
3939
GPUd() static void sort(T* begin, T* end, const S& comp);
4040
template <class T, class S>
41-
GPUd() static void sortInBlock(T* begin, T* end, const S& comp);
41+
GPUd() static void sortInBlock(int32_t nThreads, int32_t iThread, T* begin, T* end, const S& comp);
4242
template <class T, class S>
4343
GPUd() static void sortDeviceDynamic(T* begin, T* end, const S& comp);
4444
#if __cplusplus >= 202002L // sortOnDevice takes an auto parameter
@@ -268,29 +268,29 @@ GPUdi() void GPUCommonAlgorithm::sort(T* begin, T* end, const S& comp)
268268
}
269269

270270
template <class T>
271-
GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end)
271+
GPUdi() void GPUCommonAlgorithm::sortInBlock(int32_t nThreads, int32_t iThread, T* begin, T* end)
272272
{
273273
#ifndef GPUCA_GPUCODE
274274
GPUCommonAlgorithm::sort(begin, end);
275275
#else
276-
GPUCommonAlgorithm::sortInBlock(begin, end, [](auto&& x, auto&& y) { return x < y; });
276+
GPUCommonAlgorithm::sortInBlock(nThreads, iThread, begin, end, [](auto&& x, auto&& y) { return x < y; });
277277
#endif
278278
}
279279

280280
template <class T, class S>
281-
GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end, const S& comp)
281+
GPUdi() void GPUCommonAlgorithm::sortInBlock(int32_t nThreads, int32_t iThread, T* begin, T* end, const S& comp)
282282
{
283283
#ifndef GPUCA_GPUCODE
284284
GPUCommonAlgorithm::sort(begin, end, comp);
285285
#elif defined(GPUCA_DETERMINISTIC_MODE) // Not using GPUCA_DETERMINISTIC_CODE, which is enforced in TPC compression
286-
if (get_local_id(0) == 0) {
286+
if (iThread == 0) {
287287
GPUCommonAlgorithm::sort(begin, end, comp);
288288
}
289289
GPUbarrier();
290290
#else
291291
int32_t n = end - begin;
292292
for (int32_t i = 0; i < n; i++) {
293-
for (int32_t tIdx = get_local_id(0); tIdx < n; tIdx += get_local_size(0)) {
293+
for (int32_t tIdx = iThread; tIdx < n; tIdx += nThreads) {
294294
int32_t offset = i % 2;
295295
int32_t curPos = 2 * tIdx + offset;
296296
int32_t nextPos = curPos + 1;

‎GPU/Common/test/testGPUsortCUDA.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -96,12 +96,12 @@ __global__ void sortInThreadWithOperator(float* data, size_t dataLength)
9696

9797
__global__ void sortInBlock(float* data, size_t dataLength)
9898
{
99-
o2::gpu::CAAlgo::sortInBlock<float>(data, data + dataLength);
99+
o2::gpu::CAAlgo::sortInBlock<float>(blockDim.x, threadIdx.x, data, data + dataLength);
100100
}
101101

102102
__global__ void sortInBlockWithOperator(float* data, size_t dataLength)
103103
{
104-
o2::gpu::CAAlgo::sortInBlock(data, data + dataLength, [](float a, float b) { return a < b; });
104+
o2::gpu::CAAlgo::sortInBlock(blockDim.x, threadIdx.x, data, data + dataLength, [](float a, float b) { return a < b; });
105105
}
106106
///////////////////////////////////////////////////////////////
107107

‎GPU/GPUTracking/Base/GPUGeneralKernels.cxx‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -19,21 +19,21 @@ using namespace o2::gpu;
1919
template <>
2020
GPUdii() void GPUMemClean16::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors, GPUglobalref() void* ptr, uint64_t size)
2121
{
22-
const uint64_t stride = get_global_size(0);
22+
const uint64_t stride = (nBlocks * nThreads);
2323
int4 i0;
2424
i0.x = i0.y = i0.z = i0.w = 0;
2525
int4* ptra = (int4*)ptr;
2626
uint64_t len = (size + sizeof(int4) - 1) / sizeof(int4);
27-
for (uint64_t i = get_global_id(0); i < len; i += stride) {
27+
for (uint64_t i = (iBlock * nThreads + iThread); i < len; i += stride) {
2828
ptra[i] = i0;
2929
}
3030
}
3131

3232
template <>
3333
GPUdii() void GPUitoa::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors, GPUglobalref() int32_t* ptr, uint64_t size)
3434
{
35-
const uint64_t stride = get_global_size(0);
36-
for (uint64_t i = get_global_id(0); i < size; i += stride) {
35+
const uint64_t stride = (nBlocks * nThreads);
36+
for (uint64_t i = (iBlock * nThreads + iThread); i < size; i += stride) {
3737
ptr[i] = i;
3838
}
3939
}

‎GPU/GPUTracking/Base/GPUReconstructionThreading.h‎

Lines changed: 14 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -35,25 +35,25 @@ struct GPUReconstructionThreading {
3535

3636
#endif
3737

38-
#define GPUCA_TBB_KERNEL_LOOP_HOST(rec, vartype, varname, iEnd, code) \
39-
for (vartype varname = get_global_id(0); varname < iEnd; varname += get_global_size(0)) { \
40-
code \
38+
#define GPUCA_TBB_KERNEL_LOOP_HOST(rec, nBlocks, nThreads, iBlock, iThread, vartype, varname, iEnd, code) \
39+
for (vartype varname = (iBlock) * (nThreads) + (iThread); varname < iEnd; varname += (nBlocks) * (nThreads)) { \
40+
code \
4141
}
4242

4343
#ifdef GPUCA_GPUCODE
4444
#define GPUCA_TBB_KERNEL_LOOP GPUCA_TBB_KERNEL_LOOP_HOST
4545
#else
46-
#define GPUCA_TBB_KERNEL_LOOP(rec, vartype, varname, iEnd, code) \
47-
if (!rec.GetProcessingSettings().inKernelParallel) { \
48-
rec.mThreading->activeThreads->execute([&] { \
49-
tbb::parallel_for(tbb::blocked_range<vartype>(get_global_id(0), iEnd, get_global_size(0)), [&](const tbb::blocked_range<vartype>& _r_internal) { \
50-
for (vartype varname = _r_internal.begin(); varname < _r_internal.end(); varname += get_global_size(0)) { \
51-
code \
52-
} \
53-
}); \
54-
}); \
55-
} else { \
56-
GPUCA_TBB_KERNEL_LOOP_HOST(rec, vartype, varname, iEnd, code) \
46+
#define GPUCA_TBB_KERNEL_LOOP(rec, nBlocks, nThreads, iBlock, iThread, vartype, varname, iEnd, code) \
47+
if (!rec.GetProcessingSettings().inKernelParallel) { \
48+
rec.mThreading->activeThreads->execute([&] { \
49+
tbb::parallel_for(tbb::blocked_range<vartype>((iBlock) * (nThreads) + (iThread), iEnd, (nBlocks) * (nThreads)), [&](const tbb::blocked_range<vartype>& _r_internal) { \
50+
for (vartype varname = _r_internal.begin(); varname < _r_internal.end(); varname += (nBlocks) * (nThreads)) { \
51+
code \
52+
} \
53+
}); \
54+
}); \
55+
} else { \
56+
GPUCA_TBB_KERNEL_LOOP_HOST(rec, nBlocks, nThreads, iBlock, iThread, vartype, varname, iEnd, code) \
5757
}
5858
#endif
5959

‎GPU/GPUTracking/Base/hip/test/testGPUsortHIP.hip‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -104,12 +104,12 @@ __global__ void sortInThreadWithOperator(float* data, size_t dataLength)
104104

105105
__global__ void sortInBlock(float* data, size_t dataLength)
106106
{
107-
o2::gpu::CAAlgo::sortInBlock<float>(data, data + dataLength);
107+
o2::gpu::CAAlgo::sortInBlock<float>(blockDim.x, threadIdx.x, data, data + dataLength);
108108
}
109109

110110
__global__ void sortInBlockWithOperator(float* data, size_t dataLength)
111111
{
112-
o2::gpu::CAAlgo::sortInBlock(data, data + dataLength, [](float a, float b) { return a < b; });
112+
o2::gpu::CAAlgo::sortInBlock(blockDim.x, threadIdx.x, data, data + dataLength, [](float a, float b) { return a < b; });
113113
}
114114
///////////////////////////////////////////////////////////////
115115

‎GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,7 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step0at
3333
const GPUParam& GPUrestrict() param = processors.param;
3434

3535
int32_t myTrack = 0;
36-
for (uint32_t i = get_global_id(0); i < ioPtrs.nMergedTracks; i += get_global_size(0)) {
36+
for (uint32_t i = (iBlock * nThreads + iThread); i < ioPtrs.nMergedTracks; i += (nBlocks * nThreads)) {
3737
GPUbarrierWarp();
3838
const GPUTPCGMMergedTrack& GPUrestrict() trk = ioPtrs.mergedTracks[i];
3939
if (!trk.OK()) {
@@ -274,22 +274,22 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step1un
274274
static_assert(GPUCA_GET_THREAD_COUNT(GPUCA_LB_GPUTPCCompressionKernels_step1unattached) * 2 <= constants::TPC_COMP_CHUNK_SIZE);
275275
#endif
276276
#ifdef GPUCA_DETERMINISTIC_MODE
277-
CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZPadTime>(clusters->clusters[iSector][iRow]));
277+
CAAlgo::sortInBlock(nThreads, iThread, sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZPadTime>(clusters->clusters[iSector][iRow]));
278278
#else // GPUCA_DETERMINISTIC_MODE
279279
if (param.rec.tpc.compressionSortOrder == GPUSettings::SortZPadTime) {
280-
CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZPadTime>(clusters->clusters[iSector][iRow]));
280+
CAAlgo::sortInBlock(nThreads, iThread, sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZPadTime>(clusters->clusters[iSector][iRow]));
281281
} else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortZTimePad) {
282-
CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZTimePad>(clusters->clusters[iSector][iRow]));
282+
CAAlgo::sortInBlock(nThreads, iThread, sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZTimePad>(clusters->clusters[iSector][iRow]));
283283
} else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortPad) {
284-
CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortPad>(clusters->clusters[iSector][iRow]));
284+
CAAlgo::sortInBlock(nThreads, iThread, sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortPad>(clusters->clusters[iSector][iRow]));
285285
} else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortTime) {
286-
CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortTime>(clusters->clusters[iSector][iRow]));
286+
CAAlgo::sortInBlock(nThreads, iThread, sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortTime>(clusters->clusters[iSector][iRow]));
287287
}
288288
#endif // GPUCA_DETERMINISTIC_MODE
289289
GPUbarrier();
290290
}
291291

292-
for (uint32_t j = get_local_id(0); j < count; j += get_local_size(0)) {
292+
for (uint32_t j = iThread; j < count; j += nThreads) {
293293
int32_t outidx = idOffsetOut + totalCount + j;
294294
const ClusterNative& GPUrestrict() orgCl = clusters -> clusters[iSector][iRow][sortBuffer[j]];
295295

‎GPU/GPUTracking/DataCompression/GPUTPCDecompressionKernels.cxx‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -31,7 +31,7 @@ GPUdii() void GPUTPCDecompressionKernels::Thread<GPUTPCDecompressionKernels::ste
3131

3232
const uint32_t maxTime = (param.continuousMaxTimeBin + 1) * ClusterNative::scaleTimePacked - 1;
3333

34-
for (int32_t i = trackStart + get_global_id(0); i < trackEnd; i += get_global_size(0)) {
34+
for (int32_t i = trackStart + (iBlock * nThreads + iThread); i < trackEnd; i += (nBlocks * nThreads)) {
3535
uint32_t offset = decompressor.mAttachedClustersOffsets[i];
3636
TPCClusterDecompressionCore::decompressTrack(cmprClusters, param, maxTime, i, offset, decompressor);
3737
}
@@ -45,7 +45,7 @@ GPUdii() void GPUTPCDecompressionKernels::Thread<GPUTPCDecompressionKernels::ste
4545
ClusterNative* GPUrestrict() clusterBuffer = decompressor.mNativeClustersBuffer;
4646
const ClusterNativeAccess* outputAccess = decompressor.mClusterNativeAccess;
4747
uint32_t* offsets = decompressor.mUnattachedClustersOffsets;
48-
for (uint32_t i = get_global_id(0); i < GPUTPCGeometry::NROWS * nSectors; i += get_global_size(0)) {
48+
for (uint32_t i = (iBlock * nThreads + iThread); i < GPUTPCGeometry::NROWS * nSectors; i += (nBlocks * nThreads)) {
4949
uint32_t iRow = i % GPUTPCGeometry::NROWS;
5050
uint32_t iSector = sectorStart + (i / GPUTPCGeometry::NROWS);
5151
const uint32_t linearIndex = iSector * GPUTPCGeometry::NROWS + iRow;
@@ -105,7 +105,7 @@ GPUdii() void GPUTPCDecompressionUtilKernels::Thread<GPUTPCDecompressionUtilKern
105105
const GPUParam& GPUrestrict() param = processors.param;
106106
GPUTPCDecompression& GPUrestrict() decompressor = processors.tpcDecompressor;
107107
const ClusterNativeAccess* clusterAccess = decompressor.mClusterNativeAccess;
108-
for (uint32_t i = get_global_id(0); i < GPUTPCGeometry::NSECTORS * GPUTPCGeometry::NROWS; i += get_global_size(0)) {
108+
for (uint32_t i = (iBlock * nThreads + iThread); i < GPUTPCGeometry::NSECTORS * GPUTPCGeometry::NROWS; i += (nBlocks * nThreads)) {
109109
uint32_t sector = i / GPUTPCGeometry::NROWS;
110110
uint32_t row = i % GPUTPCGeometry::NROWS;
111111
for (uint32_t k = 0; k < clusterAccess->nClusters[sector][row]; k++) {
@@ -125,7 +125,7 @@ GPUdii() void GPUTPCDecompressionUtilKernels::Thread<GPUTPCDecompressionUtilKern
125125
ClusterNative* GPUrestrict() clusterBuffer = decompressor.mNativeClustersBuffer;
126126
const ClusterNativeAccess* clusterAccess = decompressor.mClusterNativeAccess;
127127
const ClusterNativeAccess* outputAccess = processors.ioPtrs.clustersNative;
128-
for (uint32_t i = get_global_id(0); i < GPUTPCGeometry::NSECTORS * GPUTPCGeometry::NROWS; i += get_global_size(0)) {
128+
for (uint32_t i = (iBlock * nThreads + iThread); i < GPUTPCGeometry::NSECTORS * GPUTPCGeometry::NROWS; i += (nBlocks * nThreads)) {
129129
uint32_t sector = i / GPUTPCGeometry::NROWS;
130130
uint32_t row = i % GPUTPCGeometry::NROWS;
131131
uint32_t count = 0;
@@ -144,7 +144,7 @@ GPUdii() void GPUTPCDecompressionUtilKernels::Thread<GPUTPCDecompressionUtilKern
144144
{
145145
ClusterNative* GPUrestrict() clusterBuffer = processors.tpcDecompressor.mNativeClustersBuffer;
146146
const ClusterNativeAccess* outputAccess = processors.ioPtrs.clustersNative;
147-
for (uint32_t i = get_global_id(0); i < GPUTPCGeometry::NSECTORS * GPUTPCGeometry::NROWS; i += get_global_size(0)) {
147+
for (uint32_t i = (iBlock * nThreads + iThread); i < GPUTPCGeometry::NSECTORS * GPUTPCGeometry::NROWS; i += (nBlocks * nThreads)) {
148148
uint32_t sector = i / GPUTPCGeometry::NROWS;
149149
uint32_t row = i % GPUTPCGeometry::NROWS;
150150
ClusterNative* buffer = clusterBuffer + outputAccess->clusterOffset[sector][row];

‎GPU/GPUTracking/Definitions/GPUDef.h‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -39,19 +39,19 @@
3939
#ifdef GPUCA_GPUCODE
4040
#define GPUCA_MAKE_SHARED_REF(vartype, varname, varglobal, varshared) const GPUsharedref() vartype& __restrict__ varname = varshared;
4141
#define GPUCA_SHARED_STORAGE(storage) storage
42-
#define GPUCA_SHARED_CACHE(target, src, size) \
42+
#define GPUCA_SHARED_CACHE(nThreads, iThread, target, src, size) \
4343
static_assert((size) % sizeof(int32_t) == 0, "Invalid shared cache size"); \
44-
for (uint32_t i_shared_cache = get_local_id(0); i_shared_cache < (size) / sizeof(int32_t); i_shared_cache += get_local_size(0)) { \
44+
for (uint32_t i_shared_cache = (iThread); i_shared_cache < (size) / sizeof(int32_t); i_shared_cache += (nThreads)) { \
4545
reinterpret_cast<GPUsharedref() int32_t*>(target)[i_shared_cache] = reinterpret_cast<GPUglobalref() const int32_t*>(src)[i_shared_cache]; \
4646
}
47-
#define GPUCA_SHARED_CACHE_REF(target, src, size, reftype, ref) \
48-
GPUCA_SHARED_CACHE(target, src, size) \
47+
#define GPUCA_SHARED_CACHE_REF(nThreads, iThread, target, src, size, reftype, ref) \
48+
GPUCA_SHARED_CACHE(nThreads, iThread, target, src, size) \
4949
GPUsharedref() const reftype* __restrict__ ref = (target)
5050
#else
5151
#define GPUCA_MAKE_SHARED_REF(vartype, varname, varglobal, varshared) const GPUglobalref() vartype & __restrict__ varname = varglobal;
5252
#define GPUCA_SHARED_STORAGE(storage)
53-
#define GPUCA_SHARED_CACHE(target, src, size)
54-
#define GPUCA_SHARED_CACHE_REF(target, src, size, reftype, ref) GPUglobalref() const reftype* __restrict__ ref = src
53+
#define GPUCA_SHARED_CACHE(nThreads, iThread, target, src, size)
54+
#define GPUCA_SHARED_CACHE_REF(nThreads, iThread, target, src, size, reftype, ref) GPUglobalref() const reftype* __restrict__ ref = src
5555
#endif
5656

5757
#endif //GPUTPCDEF_H

‎GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1937,7 +1937,7 @@ GPUd() void GPUTPCGMMerger::Finalize2(int32_t nBlocks, int32_t nThreads, int32_t
19371937
GPUd() void GPUTPCGMMerger::MergeLoopersInit(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
19381938
{
19391939
const float lowPtThresh = Param().rec.tpc.rejectQPtB5 * 1.1f; // Might need to merge tracks above the threshold with parts below the rejection threshold
1940-
for (uint32_t i = get_global_id(0); i < mMemory->nMergedTracks; i += get_global_size(0)) {
1940+
for (uint32_t i = (iBlock * nThreads + iThread); i < mMemory->nMergedTracks; i += (nBlocks * nThreads)) {
19411941
const auto& trk = mMergedTracks[i];
19421942
const auto& p = trk.GetParam();
19431943
const float qptabs = CAMath::Abs(p.GetQPt());
@@ -2006,7 +2006,7 @@ GPUd() void GPUTPCGMMerger::MergeLoopersMain(int32_t nBlocks, int32_t nThreads,
20062006
}
20072007
#endif
20082008

2009-
for (uint32_t i = get_global_id(0); i < mMemory->nLooperMatchCandidates; i += get_global_size(0)) {
2009+
for (uint32_t i = (iBlock * nThreads + iThread); i < mMemory->nLooperMatchCandidates; i += (nBlocks * nThreads)) {
20102010
for (uint32_t j = i + 1; j < mMemory->nLooperMatchCandidates; j++) {
20112011
// int32_t bs = 0;
20122012
assert(CAMath::Abs(candidates[i].refz) <= CAMath::Abs(candidates[j].refz));

‎GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -22,7 +22,7 @@ template <>
2222
GPUdii() void GPUTPCGMMergerTrackFit::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int32_t mode)
2323
{
2424
const int32_t iEnd = mode == -1 ? merger.Memory()->nRetryRefit : merger.NMergedTracks();
25-
GPUCA_TBB_KERNEL_LOOP(merger.GetRec(), int32_t, ii, iEnd, {
25+
GPUCA_TBB_KERNEL_LOOP(merger.GetRec(), nBlocks, nThreads, iBlock, iThread, int32_t, ii, iEnd, {
2626
const int32_t i = mode == -1 ? merger.RetryRefitIds()[ii] : mode ? merger.TrackOrderProcess()[ii] : ii;
2727
GPUTPCGMTrackParam::RefitTrack(merger.MergedTracks()[i], i, &merger, mode == -1);
2828
});
@@ -31,7 +31,7 @@ GPUdii() void GPUTPCGMMergerTrackFit::Thread<0>(int32_t nBlocks, int32_t nThread
3131
template <>
3232
GPUdii() void GPUTPCGMMergerFollowLoopers::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger)
3333
{
34-
GPUCA_TBB_KERNEL_LOOP(merger.GetRec(), uint32_t, i, merger.Memory()->nLoopData, {
34+
GPUCA_TBB_KERNEL_LOOP(merger.GetRec(), nBlocks, nThreads, iBlock, iThread, uint32_t, i, merger.Memory()->nLoopData, {
3535
GPUTPCGMTrackParam::PropagateLooper(&merger, i);
3636
});
3737
}

0 commit comments

Comments
 (0)