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..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 @@ -286,5 +296,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/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/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/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/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) { 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}}")