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 1/2] 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) { From a5a4c0db3f086d2831b173476e9a5c34e0c86019 Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Fri, 9 Oct 2026 13:18:18 +0200 Subject: [PATCH 2/2] GPU: make the kernel entry-point signature work on Metal MSL has no by-value kernel parameters: every one must carry an attribute naming a buffer or a builtin. It also has no ambient work-item builtins, so the grid dimensions arrive as attributes on the entry point too. GPUCA_KRNLGPU_DEF therefore gets two hooks, GPUCA_KRNL_SECTOR_ARG and GPUCA_KRNL_GRID_ARGS, which the Metal source fills in with a buffer and the four grid attributes. Both default to what the signature had, so CUDA, HIP and OpenCL generate exactly the same entry point as before. The attributes are named after the nBlocks, nThreads, iBlock and iThread that Thread() already takes, so the get_*() helpers resolve at the entry point and in everything it calls. Metal still needs its own definitions of them because the host ones assume one thread per block: get_local_id() is 0 and get_local_size() is 1 there. The generated arguments need an attribute too, each with a distinct buffer index, and the preprocessor cannot supply one: the kernel list splices arguments as a flat comma-separated list, and __COUNTER__ is monotonic across the translation unit rather than per kernel. o2_gpu_add_kernel already walks the arguments in pairs, so it emits the index there. Declarations go through GPUPtr1(idx, type, name) for pointers and GPUArg1(idx, type, name) for scalars, which each backend defines as it needs. Indices start at 3, after gpu_mem, the constant memory and the sector; the sector itself is nothing special, it just lives in the fixed part of the macro rather than in the generated list. Metal masks pointers as a 64-bit address exactly as OpenCL does, and for the same reason: GPUTRDTrackerKernels takes a GPUTRDTrackerGPU*, and a pointer to a derived class is not a valid kernel argument type there either. Binding POD pointers directly would have worked but would not have covered that case, so both go the same way. On the way back in, GPUPtr2 casts through device before handing the pointer to Thread(): the kernel's own buffers are device memory, but the Thread() entry points take the pointer unannotated, which in MSL means generic. Generated entry points are byte-identical for CUDA, HIP and OpenCL. Kernel list diagnostics: 408 to 0, and the translation unit 1108 to 881. --- GPU/Common/GPUCommonDefAPI.h | 10 +++++++ .../Base/GPUReconstructionKernelMacros.h | 11 +++++++- .../Base/metal/GPUReconstructionMETAL.metal | 10 +++++++ GPU/GPUTracking/Definitions/GPUDef.h | 26 +++++++++++++------ GPU/GPUTracking/cmake/kernel_helpers.cmake | 6 +++-- 5 files changed, 52 insertions(+), 11 deletions(-) diff --git a/GPU/Common/GPUCommonDefAPI.h b/GPU/Common/GPUCommonDefAPI.h index 357d304550513..9aa88f511f889 100644 --- a/GPU/Common/GPUCommonDefAPI.h +++ b/GPU/Common/GPUCommonDefAPI.h @@ -277,6 +277,16 @@ #define get_group_id(dim) (blockIdx.x) #elif defined(__OPENCL__) // Using OpenCL defaults +#elif defined(__METAL__) + // MSL has no work-item builtins. They arrive as attributes on the entry point, + // named there after the nBlocks / nThreads / iBlock / iThread that Thread() + // already takes, so these resolve both there and in every function below it. + #define get_global_id(dim) (iBlock * nThreads + iThread) + #define get_global_size(dim) (nBlocks * nThreads) + #define get_num_groups(dim) (nBlocks) + #define get_local_id(dim) (iThread) + #define get_local_size(dim) (nThreads) + #define get_group_id(dim) (iBlock) #else #define get_global_id(dim) iBlock #define get_global_size(dim) nBlocks diff --git a/GPU/GPUTracking/Base/GPUReconstructionKernelMacros.h b/GPU/GPUTracking/Base/GPUReconstructionKernelMacros.h index cc1c62bed507d..0887cadd7338d 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionKernelMacros.h +++ b/GPU/GPUTracking/Base/GPUReconstructionKernelMacros.h @@ -63,8 +63,17 @@ #define GPUCA_ATTRRES(...) GPUCA_M_EXPAND(GPUCA_M_CAT(GPUCA_ATTRRES_, GPUCA_M_FIRST(__VA_ARGS__)))(__VA_ARGS__) // GPU Kernel entry point +// MSL requires every kernel parameter to carry an attribute, and supplies the +// grid dimensions the same way, so the backend gets to shape both ends of the +// parameter list. +#ifndef GPUCA_KRNL_SECTOR_ARG +#define GPUCA_KRNL_SECTOR_ARG int32_t _iSector_internal +#endif +#ifndef GPUCA_KRNL_GRID_ARGS +#define GPUCA_KRNL_GRID_ARGS +#endif #define GPUCA_KRNLGPU_DEF(x_class, x_attributes, x_arguments, ...) \ - GPUg() void GPUCA_ATTRRES(GPUCA_M_STRIP(x_attributes)) GPUCA_M_CAT(krnl_, GPUCA_M_KRNL_NAME(x_class))(GPUCA_CONSMEM_PTR int32_t _iSector_internal GPUCA_M_STRIP(x_arguments)) + GPUg() void GPUCA_ATTRRES(GPUCA_M_STRIP(x_attributes)) GPUCA_M_CAT(krnl_, GPUCA_M_KRNL_NAME(x_class))(GPUCA_CONSMEM_PTR GPUCA_KRNL_SECTOR_ARG GPUCA_M_STRIP(x_arguments) GPUCA_KRNL_GRID_ARGS) #ifdef GPUCA_KRNL_DEFONLY #define GPUCA_KRNLGPU(...) GPUCA_KRNLGPU_DEF(__VA_ARGS__); diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal b/GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal index f947250e491be..14f7986f606a5 100644 --- a/GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal @@ -70,6 +70,16 @@ using namespace metal; device char* pConstantRaw [[buffer(1)]], #define GPUCA_CONSMEM (*(device GPUConstantMem*)pConstantRaw) +// Every kernel parameter needs an attribute, so the sector index arrives as a +// buffer rather than by value, and the grid dimensions come in at the end, where +// GPUCommonDefAPI.h's get_group_id() and friends pick them up. +#define GPUCA_KRNL_SECTOR_ARG constant int32_t& _iSector_internal [[buffer(2)]] +#define GPUCA_KRNL_GRID_ARGS \ + , uint iBlock [[threadgroup_position_in_grid]] \ + , uint iThread [[thread_position_in_threadgroup]] \ + , uint nThreads [[threads_per_threadgroup]] \ + , uint nBlocks [[threadgroups_per_grid]] + #include "GPUReconstructionKernelList.h" // clang-format on diff --git a/GPU/GPUTracking/Definitions/GPUDef.h b/GPU/GPUTracking/Definitions/GPUDef.h index 692e0c5ebe231..54c180d5f7187 100644 --- a/GPU/GPUTracking/Definitions/GPUDef.h +++ b/GPU/GPUTracking/Definitions/GPUDef.h @@ -21,17 +21,27 @@ #include "GPUDefParametersWrapper.h" #include "GPUCommonRtypes.h" -// Macros for masking ptrs in OpenCL kernel calls as uint64_t (The API only allows us to pass buffer objects) +// Macros for kernel arguments. OpenCL can only pass buffer objects, so pointers +// are masked as uint64_t and cast back inside the kernel. MSL needs an explicit +// buffer index on every parameter, but can bind a pointer directly. The index is +// emitted per argument by o2_gpu_add_kernel; 0, 1 and 2 are taken by gpu_mem, +// the constant memory and the sector index. #ifdef __OPENCL__ - #define GPUPtr1(a, b) uint64_t b - #ifdef __OPENCL__ - #define GPUPtr2(a, b) ((__generic a) (a) b) - #else - #define GPUPtr2(a, b) ((__global a) (a) b) - #endif + #define GPUPtr1(idx, a, b) uint64_t b + #define GPUPtr2(a, b) ((__generic a) (a) b) + #define GPUArg1(idx, a, b) a b +#elif defined(__METAL__) + // As for OpenCL, pointers travel as a 64-bit address: a pointer to a derived + // class is not a valid kernel argument type in MSL either. + #define GPUPtr1(idx, a, b) constant uint64_t& b [[buffer(idx)]] + // through device and then to generic: the kernel's own buffers are device + // memory, but the Thread() entry points take the pointer unannotated + #define GPUPtr2(a, b) ((a)((device a)(b))) + #define GPUArg1(idx, a, b) constant a& b [[buffer(idx)]] #else - #define GPUPtr1(a, b) a b + #define GPUPtr1(idx, a, b) a b #define GPUPtr2(a, b) b + #define GPUArg1(idx, a, b) a b #endif #define GPUCA_EVDUMP_FILE "event" diff --git a/GPU/GPUTracking/cmake/kernel_helpers.cmake b/GPU/GPUTracking/cmake/kernel_helpers.cmake index cc50d28ecef9e..459165d86cf5d 100644 --- a/GPU/GPUTracking/cmake/kernel_helpers.cmake +++ b/GPU/GPUTracking/cmake/kernel_helpers.cmake @@ -55,11 +55,13 @@ function(o2_gpu_add_kernel kernel_name kernel_files) math(EXPR n "${n} - 1") foreach(i RANGE 3 ${n} 2) math(EXPR j "${i} + 1") + # buffer indices 0, 1 and 2 are gpu_mem, the constant memory and the sector + math(EXPR TMP_ARG_IDX "3 + (${i} - 3) / 2") if(${ARGV${i}} MATCHES "\\*$") - string(APPEND OPT1 ",GPUPtr1(${ARGV${i}},${ARGV${j}})") + string(APPEND OPT1 ",GPUPtr1(${TMP_ARG_IDX},${ARGV${i}},${ARGV${j}})") string(APPEND OPT2 ",GPUPtr2(${ARGV${i}},${ARGV${j}})") else() - string(APPEND OPT1 ",${ARGV${i}} ${ARGV${j}}") + string(APPEND OPT1 ",GPUArg1(${TMP_ARG_IDX},${ARGV${i}},${ARGV${j}})") string(APPEND OPT2 ",${ARGV${j}}") endif() string(APPEND OPT3 ",${ARGV${i}}")