GPU: do not place anything in the Metal constant address space

The MSL generic address space covers device, threadgroup and thread but not
constant, so a generic member function cannot be called on an object that lives
in constant memory: 'cannot initialize object parameter of type X with an
expression of type constant X'. The shared code is generic throughout, so
constant memory is simply not usable on this backend.

Metal therefore implies GPUCA_NO_CONSTANT_MEMORY, which already exists for the
other backends and redirects GPUconstant() to GPUglobal(). GPUconstantref() has
to follow it, exactly as the OpenCL block already arranges; the Metal block
hardcoded 'constant' and so kept handing out constant references whatever the
setting. It now falls through to the unannotated, and therefore generic,
fallback.

Macro expansions are unchanged for host, CUDA, HIP, OpenCL and cling.
Takes the Metal translation unit from 333 errors to 235, of which the
constant-versus-generic diagnostics drop from 91 to 14.

Author: ktf

GPU/Common/GPUCommonDef.h

@@ -58,7 +58,10 @@
   #define GPUCA_ALIGPUCODE // Part of GPUTracking library but not of interface
 #endif
 
-#if (defined(__CUDACC__) && defined(GPUCA_CUDA_NO_CONSTANT_MEMORY)) || (defined(__HIPCC__) && defined(GPUCA_HIP_NO_CONSTANT_MEMORY)) || (defined(__OPENCL__) && defined(GPUCA_OPENCL_NO_CONSTANT_MEMORY))
+// __METAL__ unconditionally: the MSL generic address space does not span
+// `constant`, so a generic member function cannot be called on an object living
+// there, and the shared code is generic throughout.
+#if (defined(__CUDACC__) && defined(GPUCA_CUDA_NO_CONSTANT_MEMORY)) || (defined(__HIPCC__) && defined(GPUCA_HIP_NO_CONSTANT_MEMORY)) || (defined(__OPENCL__) && defined(GPUCA_OPENCL_NO_CONSTANT_MEMORY)) || defined(__METAL__)
   #define GPUCA_NO_CONSTANT_MEMORY
 #elif (defined(__CUDACC__) || defined(__HIPCC__)) && !defined(GPUCA_GPUCODE_HOSTONLY)
   #define GPUCA_HAS_GLOBAL_SYMBOL_CONSTANT_MEM

GPU/Common/GPUCommonDefAPI.h

@@ -167,7 +167,9 @@
   #define GPUglobalref() device
   #define GPUsharedref() threadgroup
   #define GPUprivateref() thread
-  #define GPUconstantref() constant
+  #if !defined(GPUCA_NO_CONSTANT_MEMORY)
+    #define GPUconstantref() constant
+  #endif
   #define GPUconstexprref() GPUconstexpr()
   #define GPUdouble() float
   #define GPUbarrier() threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup)