GPU: Metal branches in the common definition macros (#15772)

Defines the GPUd()/GPUshared()/GPU*ref() family for __METAL__ and
__METAL_HOST__, and teaches GPUCommonDef.h about the Metal host and device
compilation passes.

The Metal backend targets MSL 4.1 and later only. That is what lets
GPUdDefault() expand to nothing: up to MSL 4.0 a member function's implicit
this is thread, which is wrong for objects living in device memory, and
pinning defaulted constructors to device made the same type unusable in
thread or threadgroup. MSL 4.1 makes an unannotated this generic, which is
the C++ semantics this codebase already assumes. The *ref() macros stay
explicit regardless: they are correct from the OpenCL port, explicit is never
slower than generic, and constant is not covered by generic pointers at all.

Inert unless __METAL__ or __METAL_HOST__ is defined.

Author: ktf

GPU/Common/GPUCommonDef.h

@@ -31,12 +31,12 @@
 #include "GPUCommonDefSettings.h"
 
 #if !defined(__CLING__) && !defined(G__ROOT) // No GPU code for ROOT
-  #if defined(__CUDACC__) || defined(__OPENCL__) || defined(__HIPCC__) || defined(__OPENCL_HOST__)
+  #if defined(__CUDACC__) || defined(__OPENCL__) || defined(__HIPCC__) || defined(__OPENCL_HOST__) || defined(__METAL_HOST__)
     #define GPUCA_GPUCODE // Compiled by GPU compiler
   #endif
 
   #if defined(GPUCA_GPUCODE)
-    #if defined(__CUDA_ARCH__) || defined(__OPENCL__) || defined(__HIP_DEVICE_COMPILE__)
+    #if defined(__CUDA_ARCH__) || defined(__OPENCL__) || defined(__HIP_DEVICE_COMPILE__) || defined(__METAL_VERSION__)
       #define GPUCA_GPUCODE_DEVICE // Executed on device
     #endif
     #if defined(__CUDACC__)
@@ -45,6 +45,8 @@
       #define GPUCA_GPUTYPE HIP
     #elif defined(__OPENCL__) || defined(__OPENCL_HOST__)
       #define GPUCA_GPUTYPE OCL
+    #elif defined(__METAL__) || defined(__METAL_HOST__)
+      #define GPUCA_GPUTYPE METAL
     #endif
   #endif
 #endif

GPU/Common/GPUCommonDefAPI.h

@@ -27,7 +27,7 @@
 //Define macros for GPU keywords. i-version defines inline functions.
 //All host-functions in GPU code are automatically inlined, to avoid duplicate symbols.
 //For non-inline host only functions, use no keyword at all!
-#if !defined(GPUCA_GPUCODE) || defined(__OPENCL_HOST__) // For host / ROOT dictionary
+#if !defined(GPUCA_GPUCODE) || defined(__OPENCL_HOST__) || defined(__METAL_HOST__) // For host / ROOT dictionary
   #define GPUd()                                    // device function
   #define GPUdDefault()                             // default (constructor / operator) device function
   #define GPUhdDefault()                            // default (constructor / operator) host device function
@@ -124,6 +124,53 @@
   #if (!defined(__OPENCL__) || !defined(GPUCA_NO_CONSTANT_MEMORY))
     #define GPUconstantref() GPUconstant()
   #endif
+#elif defined(__METAL__) //Defines for Metal Shading Language
+  // ADDRESS SPACES. This backend targets MSL 4.1 (macOS 27) and later only --
+  // see -std=metal4.1 in the CMakeLists, which fails the build on anything
+  // older rather than miscompiling quietly.
+  //
+  // That version is what makes the port tractable: up to MSL 4.0 a member
+  // function's implicit `this` is `thread`, which is wrong for us, since most
+  // objects the kernels touch live in `device` memory. Pinning defaulted
+  // constructors and operators to `device` was the 4.0 workaround, and it made
+  // the same type unusable in `thread` or `threadgroup`. In 4.1 an unannotated
+  // `this` is GENERIC and resolves to whichever address space the object is in
+  // -- the C++ semantics this codebase already assumes -- so GPUdDefault()
+  // needs nothing at all. The compiler resolves it statically in almost every
+  // case; it only falls back to a runtime branch where it cannot see through,
+  // such as argument buffers or dynamic libraries.
+  //
+  // The *ref() macros below stay explicit even so. They are already correct
+  // from the OpenCL port, an explicit annotation is never slower than a generic
+  // one, and `constant` is not covered by generic pointers at all.
+  #define GPUdDefault()                             // generic `this` (MSL 4.1+)
+  #define GPUd()
+  #define GPUhdDefault()
+  #define GPUdi() inline
+  #define GPUdii() inline
+  #define GPUdni()
+  #define GPUdnii()
+  #define GPUh() inline
+  #define GPUhi() inline
+  #define GPUhd() inline
+  #define GPUhdi() inline
+  #define GPUhdni()
+  #define GPUg() kernel
+  #define GPUshared() threadgroup
+  #define GPUglobal() device
+  #define GPUconstant() constant // TODO: possibly add const __restrict where possible later!
+  #define GPUconstexpr() constant
+  #define GPUprivate() thread
+  #define GPUgeneric()
+  #define GPUglobalref() device
+  #define GPUsharedref() threadgroup
+  #define GPUprivateref() thread
+  #define GPUconstantref() constant
+  #define GPUconstexprref() GPUconstexpr()
+  #define GPUdouble() float
+  #define GPUbarrier() threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup)
+  #define GPUbarrierWarp() simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup)
+  #define GPUAtomic(type) atomic<type>                      // atomic variable type
 #elif defined(__HIPCC__) //Defines for HIP
   #define GPUd() __device__
   #define GPUdDefault() __device__
@@ -230,5 +277,5 @@
   #define get_group_id(dim) iBlock
 #endif
 
-    // clang-format on
+// clang-format on
 #endif