From 4efc17b8b647a06389632ead5aa68dfbfe0a7e9a Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Mon, 21 Sep 2026 11:01:45 +0200 Subject: [PATCH 1/2] GPU: stop considering OpenCL on macOS macOS ships OpenCL 1.2, below the 2.x the OpenCL backend requires, so find_package(OpenCL) there could never produce a usable backend: the version check dropped it again a few lines later. Skip the lookup on Apple instead. With that, CUDA_ENABLED, OPENCL_ENABLED and HIP_ENABLED are all necessarily off on macOS, which makes the Darwin arm of the backend dispatch dead code. Drop it along with its warning and unindent the rest. --- GPU/GPUTracking/CMakeLists.txt | 64 ++++++++++++++++------------------ dependencies/FindO2GPU.cmake | 7 ++-- 2 files changed, 35 insertions(+), 36 deletions(-) diff --git a/GPU/GPUTracking/CMakeLists.txt b/GPU/GPUTracking/CMakeLists.txt index ca58d91212084..ce3cbca5bb197 100644 --- a/GPU/GPUTracking/CMakeLists.txt +++ b/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() diff --git a/dependencies/FindO2GPU.cmake b/dependencies/FindO2GPU.cmake index d2f426c448e12..5804b046c915e 100644 --- a/dependencies/FindO2GPU.cmake +++ b/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") From 99542dd94d0f86e082611b42815083c3cf791fb0 Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Mon, 21 Sep 2026 11:03:38 +0200 Subject: [PATCH 2/2] 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. --- GPU/Common/GPUCommonDef.h | 5 ++++- GPU/Common/GPUCommonDefAPI.h | 4 +++- 2 files changed, 7 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..7346f527e00b5 100644 --- a/GPU/Common/GPUCommonDefAPI.h +++ b/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)