From 89f314b4378853a7495cf864928a3525995a467e Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Thu, 17 Sep 2026 16:59:09 +0200 Subject: [PATCH 1/2] GPU: register the Metal reconstruction library Adds sLibMETAL and the DeviceType::METAL dispatch, plus the METAL_ENABLED cmakedefine. The loader is constructed unconditionally, as the others are, and only resolves its symbols when that device type is actually requested -- so this builds and links without the backend present. --- GPU/GPUTracking/Base/GPUReconstruction.h | 2 +- .../Base/GPUReconstructionAvailableBackends.template.h | 1 + GPU/GPUTracking/Base/GPUReconstructionLibrary.cxx | 5 +++++ 3 files changed, 7 insertions(+), 1 deletion(-) diff --git a/GPU/GPUTracking/Base/GPUReconstruction.h b/GPU/GPUTracking/Base/GPUReconstruction.h index 565836a26d233..993db1f381f17 100644 --- a/GPU/GPUTracking/Base/GPUReconstruction.h +++ b/GPU/GPUTracking/Base/GPUReconstruction.h @@ -427,7 +427,7 @@ class GPUReconstruction void* mGPULib; void* mGPUEntry; }; - static std::shared_ptr sLibCUDA, sLibHIP, sLibOCL; + static std::shared_ptr sLibCUDA, sLibHIP, sLibOCL, sLibMETAL; // Debugging struct debugInternal; diff --git a/GPU/GPUTracking/Base/GPUReconstructionAvailableBackends.template.h b/GPU/GPUTracking/Base/GPUReconstructionAvailableBackends.template.h index aaf5f23b8d855..661a39e99b20f 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionAvailableBackends.template.h +++ b/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 diff --git a/GPU/GPUTracking/Base/GPUReconstructionLibrary.cxx b/GPU/GPUTracking/Base/GPUReconstructionLibrary.cxx index 2e22d4c07e77e..af79df2a09812 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionLibrary.cxx +++ b/GPU/GPUTracking/Base/GPUReconstructionLibrary.cxx @@ -101,6 +101,10 @@ std::shared_ptr* 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::sLibCUDA(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingCUDA" LIBRARY_EXTENSION, "GPUReconstruction_Create_CUDA")); std::shared_ptr GPUReconstruction::sLibHIP(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingHIP" LIBRARY_EXTENSION, "GPUReconstruction_Create_HIP")); std::shared_ptr GPUReconstruction::sLibOCL(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingOCL" LIBRARY_EXTENSION, "GPUReconstruction_Create_OCL")); +std::shared_ptr 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) {} From 817fbbb2802e3bb3ed23a11571a7fcdda756c8f3 Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Thu, 17 Sep 2026 17:03:08 +0200 Subject: [PATCH 2/2] GPU: Apple Metal backend, off by default The backend itself: the Objective-C++ host side, the .metal kernel source and its build rules, plus the CMake to enable them. Off unless asked for. FindO2GPU.cmake leaves ENABLE_METAL=OFF on macOS and the subdirectory is gated on METAL_ENABLED, so macOS keeps running on the CPU until the whole chain is validated. An earlier draft forced AUTO unconditionally, which both enabled the backend on every Mac and silently overrode an explicit -DENABLE_METAL=OFF. Apple toolchain only. The source goes .metal -> AIR through xcrun metal and nothing else; the draft's clang --target=vulkan1.4 SPIR-V route is removed rather than left commented, as it would pull in the LLVM-SPIRV translator and a second frontend Apple neither ships nor supports. Requires -std=metal4.1, stated in METAL_FLAGS. This is not merely a minimum: MSL 4.0 resolves an unannotated this as thread rather than generic, so an older toolchain would not fail, it would build the wrong thing. MSL 4.1 targets macOS 27, and Xcode 26.6 stops at metal4.0. --- GPU/GPUTracking/Base/metal/CMakeLists.txt | 124 +++++ .../Base/metal/GPUReconstructionMETAL.metal | 87 ++++ .../Base/metal/GPUReconstructionMetal.h | 68 +++ .../Base/metal/GPUReconstructionMetal.mm | 427 ++++++++++++++++++ .../GPUReconstructionMetalIncludesHost.h | 61 +++ .../metal/GPUReconstructionMetalKernels.mm | 82 ++++ ...PUReconstructionMetalKernelsSpecialize.inc | 27 ++ GPU/GPUTracking/CMakeLists.txt | 12 +- dependencies/FindO2GPU.cmake | 34 +- 9 files changed, 919 insertions(+), 3 deletions(-) create mode 100644 GPU/GPUTracking/Base/metal/CMakeLists.txt create mode 100644 GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal create mode 100644 GPU/GPUTracking/Base/metal/GPUReconstructionMetal.h create mode 100644 GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm create mode 100644 GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h create mode 100644 GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm create mode 100644 GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernelsSpecialize.inc diff --git a/GPU/GPUTracking/Base/metal/CMakeLists.txt b/GPU/GPUTracking/Base/metal/CMakeLists.txt new file mode 100644 index 0000000000000..47a2a1ae8579b --- /dev/null +++ b/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$,$-D>" + "-I$,EXCLUDE,^/usr/include/?>,$-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 $) + # 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=$ ) diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal b/GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal new file mode 100644 index 0000000000000..80f03fd42ab23 --- /dev/null +++ b/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 +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 diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.h b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.h new file mode 100644 index 0000000000000..8913ac7e8491e --- /dev/null +++ b/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 +{ + public: + GPUReconstructionMetal(const GPUSettingsDeviceBackend& cfg); + ~GPUReconstructionMetal() override; + + template + 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 + int32_t AddKernel(); + + GPUReconstructionMetalInternals* mInternals; + float mOclVersion; + + template + S& getKernelObject(); + + int32_t GetMetalPrograms(); + + private: + int32_t AddKernels(); +}; + +} // namespace o2::gpu + +#endif // GPURECONSTRUCTIONMETAL_H diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm new file mode 100644 index 0000000000000..a722b2e221dfb --- /dev/null +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm @@ -0,0 +1,427 @@ +// 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. + +#include "GPUReconstructionMetal.h" +#include "GPUConstantMem.h" +#include "GPUDefParametersLoad.inc" +#include "GPUReconstructionMetalIncludesHost.h" + +#include +#include + +#include +#include +#include // _mh_execute_header + +#define GPUErrorReturn(...) \ + { \ + GPUError(__VA_ARGS__); \ + return (1); \ + } + +#include "utils/qGetLdBinarySymbols.h" +QGET_LD_BINARY_SYMBOLS(GPUReconstructionMetalCode_src); + +GPUReconstruction* GPUReconstruction_Create_METAL(const GPUSettingsDeviceBackend& cfg) { return new GPUReconstructionMetal(cfg); } + +GPUReconstructionMetal::GPUReconstructionMetal(const GPUSettingsDeviceBackend& cfg) : GPUReconstructionProcessing::KernelInterface(cfg, sizeof(GPUReconstructionDeviceBase)) +{ + if (mMaster == nullptr) { + mInternals = new GPUReconstructionMetalInternals; + *mParDevice = o2::gpu::internal::GPUDefParametersLoad(); + } + mDeviceBackendSettings->deviceType = DeviceType::METAL; +} + +GPUReconstructionMetal::~GPUReconstructionMetal() +{ + Exit(); // Make sure we destroy everything (in particular the ITS tracker) before we exit + if (mMaster == nullptr) { + delete mInternals; + } +} + +int32_t GPUReconstructionMetal::InitDevice_Runtime() +{ + // Propagate processing settings to PoCL runtime. + // Won't affect other OpenCL runtimes. + if (int nThreads = mProcessingSettings->nHostThreads; nThreads > 0) { + auto nThreadsStr = std::to_string(nThreads); + setenv("PMETAL_CPU_MAX_CU_COUNT", nThreadsStr.c_str(), 1); + } + + if (mMaster == nullptr) { + mInternals->device = MTLCreateSystemDefaultDevice(); + + int64_t deviceGlobalMem, deviceLocalMem; + MTLSize deviceMaxWorkGroup = mInternals->device.maxThreadsPerThreadgroup; + + std::string device_name = [mInternals->device.name UTF8String]; + // On Apple Silicon, treat recommended working set as an upper bound + deviceGlobalMem = mInternals->device.recommendedMaxWorkingSetSize; + deviceLocalMem = mInternals->device.maxThreadgroupMemoryLength; + if (GetProcessingSettings().debugLevel >= 2) { + GPUInfo("Using Metal device %s with properties:", device_name.c_str()); + GPUInfo("\tUnified Memory Architecture = %ld ", mInternals->device.hasUnifiedMemory); + GPUInfo("\tRecommended Max Working Set = %ld bytes", deviceGlobalMem); + GPUInfo("\tMax thread group memory = %ld bytes", deviceLocalMem); + GPUInfo("\tmaxWorkGroup = (%ld, %ld, %ld)", deviceMaxWorkGroup.width, deviceMaxWorkGroup.height, deviceMaxWorkGroup.depth); + GPUInfo(" "); + } + + mDeviceName = device_name.c_str(); + // Basically a random number for now. + mMaxBackendThreads = 1000; + + if (GetMetalPrograms()) { + return 1; + } + + if (GetProcessingSettings().debugLevel >= 2) { + GPUInfo("Metal program and kernels loaded successfully"); + } + + // We only support internal GPUs, so the ownership of the memory is always shared + mInternals->mem_gpu = [mInternals->device newBufferWithLength:mDeviceMemorySize options:MTLResourceStorageModeShared]; + if (mInternals->mem_gpu == nil) { + GPUErrorReturn("Metal Memory Allocation Error"); + } + + // We only support internal GPUs, so the ownership of the memory is always shared. + // FIXME: until I understand how to enable the gGPUConstantMemBufferSize constexpr + int32_t tmpGPUContantMemBufferSize = 100000; // gGPUConstantMemBufferSize + mInternals->mem_constant = [mInternals->device newBufferWithLength:tmpGPUContantMemBufferSize options:MTLResourceStorageModeShared]; + if (mInternals->mem_constant) { + GPUErrorReturn("Metal Constant Memory Allocation Error"); + } + + for (int32_t i = 0; i < mNStreams; i++) { + mInternals->commandQueues[i] = [mInternals->device newCommandQueue]; + if (mInternals->commandQueues[i] == nil) { + GPUErrorReturn("Error creating Metal command queue"); + } + mInternals->commandBuffers[i] = [mInternals->commandQueues[i] commandBuffer]; + if (mInternals->commandBuffers[i] == nil) { + GPUErrorReturn("Error creating Metal command buffer"); + } + } + + mInternals->mem_host = [mInternals->device newBufferWithLength:mHostMemorySize options:MTLResourceStorageModeShared]; + if (mInternals->mem_host == nil) { + GPUErrorReturn("Error allocating pinned host memory"); + } + + mHostMemoryBase = mInternals->mem_host.contents; + mHostMemorySize = mInternals->mem_host.allocatedSize; + mDeviceMemoryBase = mInternals->mem_gpu.contents; + mDeviceMemorySize = mInternals->mem_gpu.allocatedSize; + mDeviceConstantMem = (GPUConstantMem*)mInternals->mem_constant.contents; + + if (GetProcessingSettings().debugLevel >= 1) { + GPUInfo("Memory ptrs: GPU (%ld bytes): %p - Host (%ld bytes): %p", (int64_t)mDeviceMemorySize, mDeviceMemoryBase, (int64_t)mHostMemorySize, mHostMemoryBase); + memset(mHostMemoryBase, 0xDD, mHostMemorySize); + } + + GPUInfo("Metal Initialisation successfull"); + } else { + auto* master = dynamic_cast(mMaster); + mWarpSize = master->mWarpSize; + mMaxBackendThreads = master->mMaxBackendThreads; + mDeviceName = master->mDeviceName; + mDeviceConstantMem = master->mDeviceConstantMem; + mInternals = master->mInternals; + } + + for (uint32_t i = 0; i < mEvents.size(); i++) { + auto* events = (id*)mEvents[i].data(); + new (events) id[ mEvents[i].size() ]; + } + + return (0); +} + +int32_t GPUReconstructionMetal::ExitDevice_Runtime() +{ + // Uninitialize OPENCL + SynchronizeGPU(); + + if (mMaster == nullptr) { + if (mDeviceMemoryBase) { + [mInternals->mem_gpu release]; + [mInternals->mem_constant release]; + for (uint32_t i = 0; i < mInternals->functions.size(); i++) { + [mInternals->functions[i] release]; + } + mInternals->functions.clear(); + } + if (mHostMemoryBase) { + for (int32_t i = 0; i < mNStreams; i++) { + [mInternals->commandQueues[i] release]; + [mInternals->commandBuffers[i] release]; + } + [mInternals->mem_host release]; + } + + [mInternals->library release]; + [mInternals->device release]; + GPUInfo("Metal disposed correctly"); + } + mDeviceMemoryBase = nullptr; + mHostMemoryBase = nullptr; + + return (0); +} + +size_t GPUReconstructionMetal::GPUMemCpy(void* dst, const void* src, size_t sizeBytes, int32_t stream, int32_t toGPU, deviceEvent* ev, deviceEvent* evList, int32_t nEvents) +{ + if (evList == nullptr) { + nEvents = 0; + } + if (GetProcessingSettings().debugLevel >= 3) { + stream = -1; + } + + if (stream == -1) { + SynchronizeGPU(); + } + + auto realStream = stream == -1 ? 0 : stream; + id cb = mInternals->commandBuffers[realStream]; + id blit = [cb blitCommandEncoder]; + id sourceBuffer = nil; + id destBuffer = nil; + ptrdiff_t sourceOffset = 0; + ptrdiff_t destOffset = 0; + + // Sigh. + if (src > mHostMemoryBase && src < ((char*)mHostMemoryBase + mHostMemorySize)) { + sourceBuffer = mInternals->mem_host; + sourceOffset = (char*)src - (char*)mHostMemoryBase; + } else if (src > mDeviceMemoryBase && src < ((char*)mDeviceMemoryBase + mDeviceMemorySize)) { + sourceBuffer = mInternals->mem_gpu; + sourceOffset = (char*)src - (char*)mDeviceMemoryBase; + } else { + GPUErrorReturn("Unknown buffer at %x", src); + } + + if (dst > mHostMemoryBase && dst < ((char*)mHostMemoryBase + mHostMemorySize)) { + destBuffer = mInternals->mem_host; + destOffset = (char*)src - (char*)mHostMemoryBase; + } else if (dst > mDeviceMemoryBase && dst < ((char*)mDeviceMemoryBase + mDeviceMemorySize)) { + destBuffer = mInternals->mem_gpu; + destOffset = (char*)dst - (char*)mDeviceMemoryBase; + } else { + GPUErrorReturn("Unknown buffer at %x", src); + } + + [blit copyFromBuffer:sourceBuffer + sourceOffset:sourceOffset + toBuffer:destBuffer + destinationOffset:destOffset + size:sizeBytes]; + + [blit endEncoding]; + [cb commit]; + + if (GetProcessingSettings().serializeGPU & 2) { + GPUDebug(("GPUMemCpy " + std::to_string(toGPU)).c_str(), stream, true); + } + return sizeBytes; +} + +size_t GPUReconstructionMetal::WriteToConstantMemory(size_t offset, const void* src, size_t size, int32_t stream, deviceEvent* ev) +{ + if (stream == -1) { + SynchronizeGPU(); + } + + auto realStream = stream == -1 ? 0 : stream; + id cb = mInternals->commandBuffers[realStream]; + id blit = [cb blitCommandEncoder]; + id sourceBuffer = nil; + ptrdiff_t sourceOffset = 0; + if (src > mHostMemoryBase && src < ((char*)mHostMemoryBase + mHostMemorySize)) { + sourceBuffer = mInternals->mem_host; + sourceOffset = (char*)src - (char*)mHostMemoryBase; + } else if (src > mDeviceMemoryBase && src < ((char*)mDeviceMemoryBase + mDeviceMemorySize)) { + sourceBuffer = mInternals->mem_gpu; + sourceOffset = (char*)src - (char*)mDeviceMemoryBase; + } else { + GPUErrorReturn("Unknown buffer at %x", src); + } + [blit copyFromBuffer:sourceBuffer + sourceOffset:sourceOffset + toBuffer:mInternals->mem_constant + destinationOffset:offset + size:size]; + + [blit endEncoding]; + [cb commit]; + + if (GetProcessingSettings().serializeGPU & 2) { + GPUDebug("WriteToConstantMemory", stream, true); + } + return size; +} + +void GPUReconstructionMetal::ReleaseEvent(deviceEvent ev) +{ + // FIXME: is this supposed to reset the event for it to be repurposed + // or to decrease the ref count? + auto mtlEvent = (__bridge id)(ev.get()); + [mtlEvent setSignaledValue:0]; +} + +void GPUReconstructionMetal::RecordMarker(deviceEvent* ev, int32_t stream) +{ + id cb = mInternals->commandBuffers[stream]; + // Does not change the retain count, so it's important we manage + // the lifetime of the events outside here. + auto mtlEvent = (__bridge id)(ev->get()); + [cb encodeSignalEvent:mtlEvent value:1]; + [cb commit]; +} + +int32_t GPUReconstructionMetal::DoStuckProtection(int32_t stream, deviceEvent event) +{ + if (GetProcessingSettings().stuckProtection) { + GPUError("Stuck protection not implemented for Metal"); + } else { + [mInternals->commandBuffers[stream] waitUntilCompleted]; + } + return 0; +} + +void GPUReconstructionMetal::SynchronizeGPU() +{ + for (int32_t i = 0; i < mNStreams; i++) { + [mInternals->commandBuffers[i] waitUntilCompleted]; + } +} + +void GPUReconstructionMetal::SynchronizeStream(int32_t stream) +{ + [mInternals->commandBuffers[stream] waitUntilCompleted]; +} + +void GPUReconstructionMetal::SynchronizeEvents(deviceEvent* evList, int32_t nEvents) +{ + // I wait for everything to complete for now... + for (int32_t si = 0; si < mNStreams; si++) { + id cb = mInternals->commandBuffers[si]; + [cb waitUntilCompleted]; + } +} + +void GPUReconstructionMetal::StreamWaitForEvents(int32_t stream, deviceEvent* evList, int32_t nEvents) +{ + // Encode commands to wait for all the events + id cb = mInternals->commandBuffers[stream]; + for (int32_t ei = 0; ei < nEvents; ei++) { + auto mtlEvent = (__bridge id)(evList[ei].get()); + [cb encodeWaitForEvent:mtlEvent value:1]; + } + [cb commit]; + [cb waitUntilCompleted]; +} + +bool GPUReconstructionMetal::IsEventDone(deviceEvent* evList, int32_t nEvents) +{ + for (int32_t i = 0; i < nEvents; i++) { + auto mtlEvent = (__bridge id)(evList[i].get()); + if (mtlEvent.signaledValue == 0) { + return false; + } + } + return true; +} + +int32_t GPUReconstructionMetal::GPUDebug(const char* state, int32_t stream, bool force) +{ + // Wait for Metal-Kernel to finish and check for Metal errors afterwards, in case of debugmode + if (!force && GetProcessingSettings().debugLevel <= 0) { + return (0); + } + for (int32_t si = 0; si < mNStreams; si++) { + [mInternals->commandBuffers[si] waitUntilCompleted]; + } + if (GetProcessingSettings().debugLevel >= 3) { + GPUInfo("GPU Sync Done"); + } + return (0); +} + +int32_t GPUReconstructionMetal::GPUChkErrInternal(const int64_t error, const char* file, int32_t line) const +{ + // Not sure how metal returns errors. + if (error != 0) { + GPUError("Metal Error: %ld / %s (%s:%d)", error, "Unknown", file, line); + } + return error != 0; +} + +// Return pointer+size for (__DATA|__DATA_CONST, "__gpu_resource") from the image +// that matches `image_name_substr` (e.g. "libO2GPUReconstruction.dylib"). +static const uint8_t* find_gpu_resource_in_image(const char* image_name_substr, + unsigned long* out_size) +{ + uint32_t count = _dyld_image_count(); + for (uint32_t i = 0; i < count; ++i) { + const char* name = _dyld_get_image_name(i); + if (!name || !strstr(name, image_name_substr)) { + continue; + } + + const struct mach_header* mh = _dyld_get_image_header(i); + + const auto* mh64 = (const struct mach_header_64*)mh; + const auto* p = (const uint8_t*) + getsectiondata(mh64, "__DATA", "__gpu_resource", out_size); + if (!p) { + p = (const uint8_t*) + getsectiondata(mh64, "__DATA_CONST", "__gpu_resource", out_size); + } + if (p) { + return p; + } + } + return nullptr; +} + +int32_t GPUReconstructionMetal::GetMetalPrograms() +{ + // No need for now... + [[maybe_unused]] const char* metalBuildFlags = GetProcessingSettings().metalOverrideSourceBuildFlags != "" ? GetProcessingSettings().metalOverrideSourceBuildFlags.c_str() : GPUCA_M_STR(GPUCA_METAL_BUILD_FLAGS); + + GPUInfo("Compiling Metal program from sources (Platform version %s)", [mInternals->device.architecture.name cStringUsingEncoding:NSUTF8StringEncoding]); + + unsigned long sz = 0; + const char* p = (char const*)find_gpu_resource_in_image("libO2GPUTrackingMETAL.dylib", &sz); + auto source = [[NSString alloc] initWithCString:p encoding:NSUTF8StringEncoding]; + + NSError* error = nil; + MTLCompileOptions* options = [[MTLCompileOptions alloc] init]; + + // Equivalent to clCreateProgramWithSource + mInternals->library = [mInternals->device newLibraryWithSource:source + options:options + error:&error]; + + if (error != nil) { + NSLog(@"%@", error); + NSLog(@"Error dump:\n%@", [error description]); + NSLog(@"Error debug dump:\n%@", [error debugDescription]); + GPUError("Error creating Metal program from binary"); + return 1; + } + + return AddKernels(); +} diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h new file mode 100644 index 0000000000000..47cea76d87854 --- /dev/null +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h @@ -0,0 +1,61 @@ +// 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 GPURECONSTRUCTIONOMETALINCLUDESHOST_H +#define GPURECONSTRUCTIONOMETALINCLUDESHOST_H + +#define GPUCA_GPUTYPE_METAL + +#import +#ifndef __METAL_VERSION__ +// __METAL_VERSION__ is only defined in device code. +#define __METAL_HOST__ +#endif + +#import + +#include +#include +#include +#include "GPULogging.h" + +#include "GPUReconstructionMetal.h" +#include "GPUReconstructionIncludes.h" +#include "GPUCommonHelpers.h" + +using namespace o2::gpu; + +#include +#include +#include +#include + +namespace o2::gpu +{ + +struct GPUReconstructionMetalInternals { + id device; + + std::array, GPUCA_MAX_STREAMS> commandQueues; // ~ cl_command_queue[] + std::array, GPUCA_MAX_STREAMS> commandBuffers; + + std::vector> functions; // ~ cl_kernel (symbols) + std::vector> pipelines; // compiled kernels + + id mem_gpu; // ~ cl_mem (device/global) + id mem_constant; // ~ cl_mem (constant-like) + id mem_host; // ~ cl_mem (host-visible) + + id library; // ~ cl_program +}; +} // namespace o2::gpu + +#endif // GPURECONSTRUCTIONOMETALINCLUDESHOST_H diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm new file mode 100644 index 0000000000000..54adb3076b60d --- /dev/null +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm @@ -0,0 +1,82 @@ +// 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. + +#include +#include "GPUReconstructionKernelIncludes.h" +#include "GPUReconstructionMetalIncludesHost.h" + +#include "GPUReconstructionMetalKernelsSpecialize.inc" +#include "GPUReconstructionProcessingKernels.inc" + +template void GPUReconstructionProcessing::KernelInterface::runKernelVirtual(const int num, const void* args); + +template +inline void GPUReconstructionMetal::runKernelBackend(const krnlSetupTime& _xyz, const Args&... args) +{ + id function = mInternals->functions[GetKernelNum()]; + auto& kExec = _xyz.x; + auto& runRange = _xyz.y; + // auto& events = _xyz.z; + // auto& t = _xyz.t; + + NSError* error = nil; + auto pso = [mInternals->device newComputePipelineStateWithFunction:function error:&error]; + id computeEncoder = [mInternals->commandBuffers[kExec.stream] computeCommandEncoder]; + + // Map buffers and states + [computeEncoder setComputePipelineState:pso]; + [computeEncoder setBuffer:mInternals->mem_gpu offset:0 atIndex:0]; + [computeEncoder setBuffer:mInternals->mem_constant offset:0 atIndex:1]; + [computeEncoder setBuffer:mInternals->mem_host offset:0 atIndex:2]; + + MTLSize gridSize = MTLSizeMake(runRange.index, 1, 1); + + NSUInteger threadGroupSize = pso.maxTotalThreadsPerThreadgroup; + if (threadGroupSize > runRange.index) { + threadGroupSize = runRange.index; + } + + MTLSize threadgroupSize = MTLSizeMake(threadGroupSize, 1, 1); + [computeEncoder dispatchThreads:gridSize + threadsPerThreadgroup:threadgroupSize]; +} + +template +int32_t GPUReconstructionMetal::AddKernel() +{ + NSString* kname = [[NSString alloc] initWithFormat:@"krnl_%s", GetKernelName()]; + + id krnl = [mInternals->library newFunctionWithName:kname]; + if (krnl == nil) { + GPUError("Error creating Metal Kernel: %s", [kname cStringUsingEncoding:NSUTF8StringEncoding]); + return 1; + } + + mInternals->functions.emplace_back(krnl); + return 0; +} + +template +S& GPUReconstructionMetal::getKernelObject() +{ + return mInternals->functions[GetKernelNum()]; +} + +int32_t GPUReconstructionMetal::AddKernels() +{ +#define GPUCA_KRNL(x_class, ...) \ + if (AddKernel()) { \ + return 1; \ + } +#include "GPUReconstructionKernelList.h" +#undef GPUCA_KRNL + return 0; +} diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernelsSpecialize.inc b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernelsSpecialize.inc new file mode 100644 index 0000000000000..1ee192cff822f --- /dev/null +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernelsSpecialize.inc @@ -0,0 +1,27 @@ +// 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 GPUReconstructionMetalKernelsSpecialize.inc + +template <> +inline void GPUReconstructionMetal::runKernelBackend(const krnlSetupTime& _xyz, void* const& ptr, uint64_t const& size) +{ + const uint64_t offset = static_cast(ptr) - static_cast(mDeviceMemoryBase); + const uint64_t length = (size + 15ull) & ~15ull; + + id cb = mInternals->commandBuffers[_xyz.x.stream]; + id blit = [cb blitCommandEncoder]; + + [blit fillBuffer:mInternals->mem_gpu range:NSMakeRange(offset, length) value:0]; + [blit endEncoding]; + + [cb commit]; +} diff --git a/GPU/GPUTracking/CMakeLists.txt b/GPU/GPUTracking/CMakeLists.txt index ca58d91212084..3cbe100a7dc4b 100644 --- a/GPU/GPUTracking/CMakeLists.txt +++ b/GPU/GPUTracking/CMakeLists.txt @@ -465,7 +465,11 @@ 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") + if(METAL_ENABLED) + add_subdirectory(Base/metal) + else() + message(WARNING "GPU Tracking disabled on MacOS") + endif() else() make_directory(${CMAKE_CURRENT_BINARY_DIR}/genGPUArch) set(GPU_CONST_PARAM_FILES) @@ -499,6 +503,12 @@ if(CUDA_ENABLED OR OPENCL_ENABLED OR HIP_ENABLED) add_subdirectory(Base/hip) endif() endif() +elseif(CMAKE_SYSTEM_NAME MATCHES Darwin AND METAL_ENABLED) + # Gated on METAL_ENABLED, which FindO2GPU.cmake leaves OFF unless the build + # explicitly asks for -DENABLE_METAL=ON. macOS keeps running on the CPU by + # default until the Metal chain is validated; adding this subdirectory + # unconditionally would have built the backend on every Mac. + add_subdirectory(Base/metal) endif() if(ALIGPU_BUILD_TYPE STREQUAL "O2" OR ALIGPU_BUILD_TYPE STREQUAL "Standalone") diff --git a/dependencies/FindO2GPU.cmake b/dependencies/FindO2GPU.cmake index d2f426c448e12..870ddd0839635 100644 --- a/dependencies/FindO2GPU.cmake +++ b/dependencies/FindO2GPU.cmake @@ -20,14 +20,34 @@ 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) + if(NOT DEFINED ENABLE_OPENCL) + set(ENABLE_OPENCL "AUTO") + endif() +else() + # OFF, not AUTO: macOS keeps running on the CPU by default until the whole + # Metal chain is validated. AUTO would enable the backend on every Mac merely + # because the frameworks are present, which is exactly what we do not want + # while it is unproven -- and it would do so silently. + # + # `if(NOT DEFINED ...)` so an explicit -DENABLE_METAL=ON is honoured; the + # earlier draft set this unconditionally and quietly overrode whatever the + # user asked for. + # + # Note it cannot build anywhere yet in any case: the backend requires + # -std=metal4.1, which needs a newer toolchain than Xcode 26.6. + if(NOT DEFINED ENABLE_METAL) + set(ENABLE_METAL "OFF") + endif() endif() if(NOT DEFINED ENABLE_HIP) set(ENABLE_HIP "AUTO") endif() string(TOUPPER "${ENABLE_CUDA}" ENABLE_CUDA) string(TOUPPER "${ENABLE_OPENCL}" ENABLE_OPENCL) +if(APPLE) + string(TOUPPER "${ENABLE_METAL}" ENABLE_METAL) +endif() string(TOUPPER "${ENABLE_HIP}" ENABLE_HIP) if(NOT DEFINED CMAKE_BUILD_TYPE_UPPER) string(TOUPPER "${CMAKE_BUILD_TYPE}" CMAKE_BUILD_TYPE_UPPER) @@ -429,6 +449,16 @@ if(ENABLE_HIP) endif() endif() +if(ENABLE_METAL) + find_library(METAL Metal) + find_library(CF CoreFoundation) + find_library(FOUNDATION Foundation) + find_library(QUARTZ_CORE QuartzCore) + + set(METAL_ENABLED ON) + set(METAL_FRAMEWORKS ${METAL} ${CF} ${FOUNDATION} ${QUARTZ_CORE}) +endif() + # if we end up here without a FATAL, it means we have found the "O2GPU" package set(O2GPU_FOUND TRUE) if (NOT GPUCA_FINDO2GPU_CHECK_ONLY)