GPU: make the kernel entry-point signature work on Metal

Two things MSL does differently from every other backend.

It has no work-item builtins: the grid dimensions arrive as kernel attributes.
Without a branch of its own Metal fell through to the host definitions of
get_group_id() and friends, which name iBlock and nBlocks, variables that only
exist in the host-side loop.

And it requires every kernel parameter to carry an attribute, so the sector
index cannot be passed by value.

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.

Kernel list diagnostics: 408 to 96, and the translation unit 1108 to 998. The 96
left are the 48 kernels that take arguments, which still need an answer for how
Metal passes them.

Author: ktf

GPU/Common/GPUCommonDefAPI.h

@@ -277,6 +277,15 @@
   #define get_group_id(dim) (blockIdx.x)
 #elif defined(__OPENCL__)
   // Using OpenCL defaults
+#elif defined(__METAL__)
+  // MSL has no work-item builtins; these come in as kernel attributes, declared
+  // by GPUCA_KRNL_GRID_ARGS on every entry point.
+  #define get_global_id(dim) (_metalTgIg * _metalTPerTg + _metalTiTg)
+  #define get_global_size(dim) (_metalTPerTg * _metalTgPerG)
+  #define get_num_groups(dim) (_metalTgPerG)
+  #define get_local_id(dim) (_metalTiTg)
+  #define get_local_size(dim) (_metalTPerTg)
+  #define get_group_id(dim) (_metalTgIg)
 #else
   #define get_global_id(dim) iBlock
   #define get_global_size(dim) nBlocks

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__);

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 _metalTgIg   [[threadgroup_position_in_grid]] \
+  , uint _metalTiTg   [[thread_position_in_threadgroup]] \
+  , uint _metalTPerTg [[threads_per_threadgroup]] \
+  , uint _metalTgPerG [[threadgroups_per_grid]]
+
 #include "GPUReconstructionKernelList.h"
 
 // clang-format on