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

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)

GPU/GPUTracking/CMakeLists.txt

@@ -464,40 +464,36 @@ endif()
 
 # Add CMake recipes for GPU Tracking librararies
 if(CUDA_ENABLED OR OPENCL_ENABLED OR HIP_ENABLED)
-  if(CMAKE_SYSTEM_NAME MATCHES Darwin)
-    message(WARNING "GPU Tracking disabled on MacOS")
-  else()
-    make_directory(${CMAKE_CURRENT_BINARY_DIR}/genGPUArch)
-    set(GPU_CONST_PARAM_FILES)
-    foreach(GPU_ARCH ${GPU_CONST_PARAM_ARCHITECTUES})
-      set(PARAMFILE ${CMAKE_CURRENT_BINARY_DIR}/genGPUArch/gpu_const_param_${GPU_ARCH}.par)
-      add_custom_command(
-        OUTPUT ${PARAMFILE}
-        COMMAND bash -c
-                "echo -e '#define GPUCA_GPUTYPE_${GPU_ARCH}\\n#define PARAMETER_FILE \"GPUDefParametersDefaults.h\"\\ngInterpreter->AddIncludePath(\"${CMAKE_CURRENT_SOURCE_DIR}/Definitions\");\\ngInterpreter->AddIncludePath(\"${ON_THE_FLY_DIR}\");\\n.x ${CMAKE_CURRENT_SOURCE_DIR}/Standalone/tools/dumpGPUDefParam.C(\"${PARAMFILE}\")\\n.q\\n'"
-                | root -l -b > /dev/null
-        VERBATIM
-        WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/genGPUArch
-        MAIN_DEPENDENCY Standalone/tools/dumpGPUDefParam.C
-        DEPENDS ${GPU_DEFAULT_PARAMS_HEADER}
-                ${GPU_DEFAULT_PARAMS_HEADER_DEVICE}
-                ${ON_THE_FLY_DIR}/GPUDefParametersLoadPrepare.h
-                ${ON_THE_FLY_DIR}/GPUDefParametersLoad.inc
-        COMMENT "Generating GPU parameter set for architecture ${GPU_ARCH}")
-      LIST(APPEND GPU_CONST_PARAM_FILES ${PARAMFILE})
-    endforeach()
-    add_custom_target(${MODULE}_GPU_CONST_PARAM_ARCHS ALL DEPENDS ${GPU_CONST_PARAM_FILES})
-    install(FILES ${GPU_CONST_PARAM_FILES} DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/GPU/arch_param)
-
-    if(CUDA_ENABLED)
-      add_subdirectory(Base/cuda)
-    endif()
-    if(OPENCL_ENABLED)
-      add_subdirectory(Base/opencl)
-    endif()
-    if(HIP_ENABLED)
-      add_subdirectory(Base/hip)
-    endif()
+  make_directory(${CMAKE_CURRENT_BINARY_DIR}/genGPUArch)
+  set(GPU_CONST_PARAM_FILES)
+  foreach(GPU_ARCH ${GPU_CONST_PARAM_ARCHITECTUES})
+    set(PARAMFILE ${CMAKE_CURRENT_BINARY_DIR}/genGPUArch/gpu_const_param_${GPU_ARCH}.par)
+    add_custom_command(
+      OUTPUT ${PARAMFILE}
+      COMMAND bash -c
+              "echo -e '#define GPUCA_GPUTYPE_${GPU_ARCH}\\n#define PARAMETER_FILE \"GPUDefParametersDefaults.h\"\\ngInterpreter->AddIncludePath(\"${CMAKE_CURRENT_SOURCE_DIR}/Definitions\");\\ngInterpreter->AddIncludePath(\"${ON_THE_FLY_DIR}\");\\n.x ${CMAKE_CURRENT_SOURCE_DIR}/Standalone/tools/dumpGPUDefParam.C(\"${PARAMFILE}\")\\n.q\\n'"
+              | root -l -b > /dev/null
+      VERBATIM
+      WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/genGPUArch
+      MAIN_DEPENDENCY Standalone/tools/dumpGPUDefParam.C
+      DEPENDS ${GPU_DEFAULT_PARAMS_HEADER}
+              ${GPU_DEFAULT_PARAMS_HEADER_DEVICE}
+              ${ON_THE_FLY_DIR}/GPUDefParametersLoadPrepare.h
+              ${ON_THE_FLY_DIR}/GPUDefParametersLoad.inc
+      COMMENT "Generating GPU parameter set for architecture ${GPU_ARCH}")
+    LIST(APPEND GPU_CONST_PARAM_FILES ${PARAMFILE})
+  endforeach()
+  add_custom_target(${MODULE}_GPU_CONST_PARAM_ARCHS ALL DEPENDS ${GPU_CONST_PARAM_FILES})
+  install(FILES ${GPU_CONST_PARAM_FILES} DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/GPU/arch_param)
+
+  if(CUDA_ENABLED)
+    add_subdirectory(Base/cuda)
+  endif()
+  if(OPENCL_ENABLED)
+    add_subdirectory(Base/opencl)
+  endif()
+  if(HIP_ENABLED)
+    add_subdirectory(Base/hip)
   endif()
 endif()
 

dependencies/FindO2GPU.cmake

@@ -20,8 +20,11 @@ set(HIP_AMDGPUTARGET_DEFAULT_MINIMAL gfx906)
 if(NOT DEFINED ENABLE_CUDA)
   set(ENABLE_CUDA "AUTO")
 endif()
-if(NOT DEFINED ENABLE_OPENCL)
-  set(ENABLE_OPENCL "AUTO")
+if(NOT APPLE)
+  # macOS ships OpenCL 1.2 only, below the 2.x that the OpenCL backend needs.
+  if(NOT DEFINED ENABLE_OPENCL)
+    set(ENABLE_OPENCL "AUTO")
+  endif()
 endif()
 if(NOT DEFINED ENABLE_HIP)
   set(ENABLE_HIP "AUTO")