GPU: keep the constant memory block out of Metal's constant address space

MSL's 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. Metal therefore implies GPUCA_NO_CONSTANT_MEMORY, which redirects
GPUconstant() to GPUglobal(). GPUconstantref() has to follow: the Metal block
hardcoded constant and so kept handing out constant references whatever the
setting; dropping it lets it fall through to the generic fallback.

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,6 @@
   #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)