GPU: Apple Metal backend, off by default

GPU/GPUTracking/Base/GPUReconstruction.h

@@ -427,7 +427,7 @@ class GPUReconstruction
     void* mGPULib;
     void* mGPUEntry;
   };
-  static std::shared_ptr<LibraryLoader> sLibCUDA, sLibHIP, sLibOCL;
+  static std::shared_ptr<LibraryLoader> sLibCUDA, sLibHIP, sLibOCL, sLibMETAL;
 
   // Debugging
   struct debugInternal;

GPU/GPUTracking/Base/GPUReconstructionAvailableBackends.template.h

@@ -16,5 +16,6 @@
 #cmakedefine CUDA_ENABLED
 #cmakedefine HIP_ENABLED
 #cmakedefine OPENCL_ENABLED
+#cmakedefine METAL_ENABLED
 #cmakedefine GPUCA_COMPILER_VERSIONS @GPUCA_COMPILER_VERSIONS@
 // clang-format on

GPU/GPUTracking/Base/GPUReconstructionLibrary.cxx

@@ -101,6 +101,10 @@ std::shared_ptr<GPUReconstruction::LibraryLoader>* GPUReconstruction::GetLibrary
   } else if (type == DeviceType::OCL) {
 #ifdef OPENCL_ENABLED
     return &sLibOCL;
+#endif
+  } else if (type == DeviceType::METAL) {
+#ifdef METAL_ENABLED
+    return &sLibMETAL;
 #endif
   } else {
     GPUError("Error: Invalid device type %u", (uint32_t)type);
@@ -125,6 +129,7 @@ GPUReconstruction* GPUReconstruction::CreateInstance(const char* type, bool forc
 std::shared_ptr<GPUReconstruction::LibraryLoader> GPUReconstruction::sLibCUDA(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingCUDA" LIBRARY_EXTENSION, "GPUReconstruction_Create_CUDA"));
 std::shared_ptr<GPUReconstruction::LibraryLoader> GPUReconstruction::sLibHIP(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingHIP" LIBRARY_EXTENSION, "GPUReconstruction_Create_HIP"));
 std::shared_ptr<GPUReconstruction::LibraryLoader> GPUReconstruction::sLibOCL(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingOCL" LIBRARY_EXTENSION, "GPUReconstruction_Create_OCL"));
+std::shared_ptr<GPUReconstruction::LibraryLoader> GPUReconstruction::sLibMETAL(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingMETAL" LIBRARY_EXTENSION, "GPUReconstruction_Create_METAL"));
 
 GPUReconstruction::LibraryLoader::LibraryLoader(const char* lib, const char* func) : mLibName(lib), mFuncName(func), mGPULib(nullptr), mGPUEntry(nullptr) {}
 

GPU/GPUTracking/Base/metal/CMakeLists.txt

@@ -0,0 +1,124 @@
+# Copyright 2019-2020 CERN and copyright holders of ALICE O2.
+# See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
+# All rights not expressly granted are reserved.
+#
+# This software is distributed under the terms of the GNU General Public
+# License v3 (GPL Version 3), copied verbatim in the file "COPYING".
+#
+# In applying this license CERN does not waive the privileges and immunities
+# granted to it by virtue of its status as an Intergovernmental Organization
+# or submit itself to any jurisdiction.
+
+set(MODULE GPUTrackingMETAL)
+enable_language(ASM)
+
+message(STATUS "Building GPUTracking with Metal support")
+
+# convenience variables
+if(ALIGPU_BUILD_TYPE STREQUAL "Standalone")
+  set(GPUDIR ${CMAKE_SOURCE_DIR}/../)
+else()
+  set(GPUDIR ${CMAKE_SOURCE_DIR}/GPU/GPUTracking)
+endif()
+set(METAL_SRC ${GPUDIR}/Base/metal/GPUReconstructionMetal.metal)
+set(METAL_BIN ${CMAKE_CURRENT_BINARY_DIR}/GPUReconstructionMetalCode)
+
+# MSL 4.1 or later, not negotiable: the address-space macros in
+# GPUCommonDefAPI.h rely on an unannotated `this` being generic, which 4.0 and
+# earlier resolve as `thread` instead -- so an older toolchain would not fail
+# cleanly, it would build the wrong thing. Stating -std here turns that into an
+# immediate "invalid value 'metal4.1'" from the compiler.
+#
+# 4.1 targets macOS 27 (MSL specification, section 1.6.10); Xcode 26.6 stops at
+# metal4.0, so this needs a newer toolchain than the CI builders currently have.
+set(METAL_FLAGS -std=metal4.1 ${GPUCA_METAL_DENORMALS_FLAGS})
+set(METAL_DEFINES "-D$<JOIN:$<TARGET_PROPERTY:O2::GPUTracking,COMPILE_DEFINITIONS>,$<SEMICOLON>-D>"
+            "-I$<JOIN:$<FILTER:$<TARGET_PROPERTY:O2::GPUTracking,INCLUDE_DIRECTORIES>,EXCLUDE,^/usr/include/?>,$<SEMICOLON>-I>"
+            -I${CMAKE_SOURCE_DIR}/Detectors/TRD/base/src
+            -I${CMAKE_SOURCE_DIR}/Detectors/Base/src
+            -I${CMAKE_SOURCE_DIR}/DataFormats/Reconstruction/src
+            -DGPUCA_GPUCODE=1
+)
+
+set(SRCS GPUReconstructionMetal.mm GPUReconstructionMetalKernels.mm)
+set(HDRS GPUReconstructionMetal.h GPUReconstructionMetalIncludesHost.h)
+
+message("FOO ${SRCS} ${METAL_ENABLED} ${ALIGPU_BUILD_TYPE}")
+
+
+if(ALIGPU_BUILD_TYPE STREQUAL "O2")
+  o2_add_library(${MODULE}
+                 SOURCES ${SRCS}
+                 PUBLIC_LINK_LIBRARIES O2::GPUTracking
+                 TARGETVARNAME targetName)
+
+  target_link_libraries(${targetName} PUBLIC ${METAL_FRAMEWORKS})
+
+  target_compile_definitions(${targetName} PRIVATE $<TARGET_PROPERTY:O2::GPUTracking,COMPILE_DEFINITIONS>)
+  # the compile_defitions are not propagated automatically on purpose (they are
+  # declared PRIVATE) so we are not leaking them outside of the GPU**
+  # directories
+endif()
+
+if(ALIGPU_BUILD_TYPE STREQUAL "Standalone")
+  add_library(${MODULE} SHARED ${SRCS})
+  target_link_libraries(${MODULE} GPUTracking)
+  install(TARGETS ${MODULE})
+  set(targetName ${MODULE})
+endif()
+
+if(METAL_ENABLED) # BUILD Metal source code for runtime compilation target
+
+  # Apple toolchain only, deliberately: the source goes .metal -> AIR with
+  # `xcrun metal` and nothing else. An earlier draft routed C++ through
+  # clang --target=vulkan1.4 to SPIR-V and then into Metal; that is not used and
+  # is not wanted -- it drags in the LLVM-SPIRV translator, a second frontend
+  # with its own dialect quirks, and a translation step Apple neither ships nor
+  # supports. Kept out rather than commented out so nobody revives it by
+  # accident.
+
+  message("Metal was enabled ${METAL_SRC}")
+  # executes clang to preprocess
+  add_custom_command(
+      OUTPUT ${METAL_BIN}.metal
+      COMMAND xcrun -sdk macosx metal
+              -Wno-unused-command-line-argument
+              ${METAL_FLAGS}
+              ${METAL_DEFINES}
+              -MD -MT ${METAL_BIN}.src -MF ${METAL_BIN}.src.d
+              -E -P ${METAL_SRC} > ${METAL_BIN}.metal
+      DEPENDS ${METAL_SRC}
+      DEPFILE ${METAL_BIN}.src.d
+      COMMAND_EXPAND_LISTS
+      COMMENT "Preparing Metal source file for run time compilation ${METAL_BIN}.metal")
+
+  # Create the ir
+  add_custom_command(
+      OUTPUT ${METAL_BIN}.ir
+      COMMAND xcrun -sdk macosx metal
+              -Wno-unused-command-line-argument
+              -Wno-c++17-extensions
+              -ferror-limit=10000
+              ${METAL_FLAGS}
+              ${METAL_DEFINES}
+              ${METAL_BIN}.metal
+              -o ${METAL_BIN}.ir
+      DEPENDS ${METAL_BIN}.metal
+      COMMAND_EXPAND_LISTS
+      COMMENT "Preparing Metal intermediate representation for run time compilation ${METAL_BIN}.ir")
+
+  add_custom_target(metal_preprocessed_code ALL DEPENDS ${METAL_BIN}.metal COMMENT "Needed to inject dependency on its creation")
+  add_custom_target(metal_intermediate_representation ALL DEPENDS ${METAL_BIN}.ir COMMENT "Needed to inject dependency on its creation")
+
+  # Pack the file into __DATA,__gpu_resource during final link. This
+  # way we do not need to create an intermediate object.
+  target_link_options(${targetName}
+                      PRIVATE
+                      "-Wl,-sectcreate,__DATA,__gpu_resource,${METAL_BIN}.metal")
+  add_dependencies(${targetName} metal_preprocessed_code)
+  add_dependencies(${targetName} metal_intermediate_representation)
+endif()
+
+install(FILES ${HDRS} DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/GPU)
+
+target_compile_definitions(${targetName} PRIVATE GPUCA_METAL_BUILD_FLAGS=$<JOIN:${METAL_FLAGS},\ > )

GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal

@@ -0,0 +1,87 @@
+// Copyright 2019-2025 CERN and copyright holders of ALICE O2.
+// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
+// All rights not expressly granted are reserved.
+//
+// This software is distributed under the terms of the GNU General Public
+// License v3 (GPL Version 3), copied verbatim in the file "COPYING".
+//
+// In applying this license CERN does not waive the privileges and immunities
+// granted to it by virtue of its status as an Intergovernmental Organization
+// or submit itself to any jurisdiction.
+
+/// \file GPUReconstructionMetal.metal
+
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wgnu-zero-variadic-macro-arguments"
+// clang-format off
+
+// --- Backend selection -------------------------------------------------------
+#define GPUCA_GPUTYPE_METAL 1
+
+// --- Metal stdlib ------------------------------------------------------------
+#include <metal_stdlib>
+using namespace metal;
+
+// --- OpenCL compatibility shims ---------------------------------------------
+
+// Address space aliases (match OpenCL vernacular used by the project)
+#define global   device
+#define local    threadgroup
+#define constant constant
+
+#ifndef M_PI
+#define M_PI 3.1415926535f
+#endif
+
+// Disable assertions inside GPU code (same as OpenCL variant)
+#ifdef assert
+# undef assert
+#endif
+#define assert(param)
+
+// --- Project headers ---------------------------------------------------------
+#if false
+#include "GPUCommonDef.h"
+#include "GPUCommonTypeTraits.h" // (MSL can't include system C headers inside kernels; these should be GPU-safe)
+#include "GPUCommonArray.h"
+#include "GPUConstantMem.h"
+// FIXME: We need a solution for the generic memory
+#include "GPUReconstructionIncludesDeviceAll.h"
+#endif
+
+// --- Kernel list expansion ---------------------------------------------------
+#define GPUCA_KRNL(...) GPUCA_KRNLGPU(__VA_ARGS__)
+
+// --- Constant memory + global heap plumbing ---------------------------------
+// In OpenCL, the kernels used:
+//   GPUglobal() char *gpu_mem, GPUconstant() GPUConstantMem* pConstant,
+// For Metal we bind them to buffer(0) and buffer(1) respectively.
+// NOTE: Metal prefers references for constant buffers; keep a reference here.
+#define GPUCA_CONSMEM_PTR \
+  device char* gpu_mem               [[buffer(0)]], \
+  constant GPUConstantMem& pConstant [[buffer(1)]],
+#define GPUCA_CONSMEM (pConstant)
+
+// If your code uses barriers like barrier(CLK_LOCAL_MEM_FENCE) via macros,
+// you likely already map them in GPUReconstructionIncludesDeviceAll.h for each backend.
+// If not, uncomment the following generic mapping:
+// #define barrier(flags) threadgroup_barrier(mem_flags::mem_threadgroup)
+
+// Include the actual kernels
+// FIXME: disabled for now. We need to find a sustainable solution to 
+// the missing __generic in Metal.
+#if 0
+#include "GPUReconstructionKernelList.h"
+#endif
+
+// Clean up local macro namespace if desired
+// #undef GPUCA_KRNL
+// #undef GPUCA_CONSMEM_PTR
+// #undef GPUCA_CONSMEM
+// #undef global
+// #undef local
+// #undef constant
+// #undef private
+
+// clang-format on
+#pragma clang diagnostic pop

GPU/GPUTracking/Base/metal/GPUReconstructionMetal.h

@@ -0,0 +1,68 @@
+// Copyright 2019-2025 CERN and copyright holders of ALICE O2.
+// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
+// All rights not expressly granted are reserved.
+//
+// This software is distributed under the terms of the GNU General Public
+// License v3 (GPL Version 3), copied verbatim in the file "COPYING".
+//
+// In applying this license CERN does not waive the privileges and immunities
+// granted to it by virtue of its status as an Intergovernmental Organization
+// or submit itself to any jurisdiction.
+
+#ifndef GPURECONSTRUCTIONMETAL_H
+#define GPURECONSTRUCTIONMETAL_H
+
+#include "GPUReconstructionDeviceBase.h"
+
+extern "C" o2::gpu::GPUReconstruction* GPUReconstruction_Create_METAL(const o2::gpu::GPUSettingsDeviceBackend& cfg);
+
+namespace o2::gpu
+{
+struct GPUReconstructionMetalInternals;
+
+class GPUReconstructionMetal : public GPUReconstructionProcessing::KernelInterface<GPUReconstructionMetal, GPUReconstructionDeviceBase>
+{
+ public:
+  GPUReconstructionMetal(const GPUSettingsDeviceBackend& cfg);
+  ~GPUReconstructionMetal() override;
+
+  template <class T, int32_t I = 0, typename... Args>
+  void runKernelBackend(const krnlSetupTime& _xyz, const Args&... args);
+
+ protected:
+  int32_t InitDevice_Runtime() override;
+  int32_t ExitDevice_Runtime() override;
+
+  virtual int32_t GPUChkErrInternal(const int64_t error, const char* file, int32_t line) const override;
+
+  void SynchronizeGPU() override;
+  int32_t DoStuckProtection(int32_t stream, deviceEvent event) override;
+  int32_t GPUDebug(const char* state = "UNKNOWN", int32_t stream = -1, bool force = false) override;
+  void SynchronizeStream(int32_t stream) override;
+  void SynchronizeEvents(deviceEvent* evList, int32_t nEvents = 1) override;
+  void StreamWaitForEvents(int32_t stream, deviceEvent* evList, int32_t nEvents = 1) override;
+  bool IsEventDone(deviceEvent* evList, int32_t nEvents = 1) override;
+
+  size_t WriteToConstantMemory(size_t offset, const void* src, size_t size, int32_t stream = -1, deviceEvent* ev = nullptr) override;
+  size_t GPUMemCpy(void* dst, const void* src, size_t size, int32_t stream, int32_t toGPU, deviceEvent* ev = nullptr, deviceEvent* evList = nullptr, int32_t nEvents = 1) override;
+  void ReleaseEvent(deviceEvent ev) override;
+  void RecordMarker(deviceEvent* ev, int32_t stream) override;
+
+  template <class T, int32_t I = 0>
+  int32_t AddKernel();
+
+  GPUReconstructionMetalInternals* mInternals;
+  float mOclVersion;
+
+  template <class S, class T, int32_t I>
+  S& getKernelObject();
+
+  int32_t GetMetalPrograms();
+
+ private:
+  int32_t AddKernels();
+};
+
+} // namespace o2::gpu
+
+#endif // GPURECONSTRUCTIONMETAL_H