From 969b395198d660d3c63ff400df3352eb6774c22d Mon Sep 17 00:00:00 2001
From: Giulio Eulisse <10544+ktf@users.noreply.github.com>
Date: Wed, 23 Sep 2026 14:18:21 +0200
Subject: [PATCH] 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.
---
 GPU/Common/GPUCommonDef.h    | 5 ++++-
 GPU/Common/GPUCommonDefAPI.h | 1 -
 2 files changed, 4 insertions(+), 2 deletions(-)

diff --git a/GPU/Common/GPUCommonDef.h b/GPU/Common/GPUCommonDef.h
index 6e2269650f576..90746019d9a99 100644
--- a/GPU/Common/GPUCommonDef.h
+++ b/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
diff --git a/GPU/Common/GPUCommonDefAPI.h b/GPU/Common/GPUCommonDefAPI.h
index a04934304d525..4cc2c8e69074d 100644
--- a/GPU/Common/GPUCommonDefAPI.h
+++ b/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)