From 691bad3044b164705ec993cadc96b09e9036feb8 Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Mon, 28 Sep 2026 20:16:56 +0200 Subject: [PATCH] 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. macOS had no usable GPU backend before this. It ships OpenCL 1.2, below the 2.x the OpenCL backend requires, so find_package(OpenCL) there could never produce one; the version check dropped it again a few lines later. That lookup is now skipped on Apple, which leaves CUDA_ENABLED, OPENCL_ENABLED and HIP_ENABLED all necessarily off, so the Darwin arm of the backend dispatch was dead code and goes with it. Metal takes its place. The build rules also give Metal its entry in the no-fast-math table, so that once GPUCA_DETERMINISTIC_MODE reaches GPUCA_DETERMINISTIC_MODE_MAP_NO_FAST_MATH it drops its fast math flags like the other backends rather than keeping whatever the build type gave it. 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. Apple toolchain only: the source goes .metal -> AIR through xcrun metal and nothing else, with no SPIR-V translation step in between. Requires -std=metal4.1, the first MSL version with a generic address space. Earlier versions reject an unannotated pointer with 'pointer type must have explicit address space qualifier' and an unannotated 'this' with 'cannot initialize object parameter', both of which GPUCommonDefAPI.h relies on for GPUgeneric() and GPUdDefault(). Verified against Xcode 27, which ships metal4.1; Xcode 26 and earlier stop at metal4.0. The Metal frameworks ship with every macOS, so finding them says nothing about whether the backend can be built; the configure compiles a three-line kernel as MSL 4.1 to answer that directly. The deployment target has no say either: -std= is what picks the target OS, and MACOSX_DEPLOYMENT_TARGET and -mmacosx-version-min are both ignored by the Metal compiler. AUTO therefore turns Metal off on an older toolchain instead of failing somewhere in the middle of the build, and an explicit ENABLE_METAL=ON says why it cannot be honoured. --- GPU/GPUTracking/Base/metal/CMakeLists.txt | 110 +++++ .../Base/metal/GPUReconstructionMETAL.metal | 77 ++++ .../Base/metal/GPUReconstructionMetal.h | 66 +++ .../Base/metal/GPUReconstructionMetal.mm | 405 ++++++++++++++++++ .../GPUReconstructionMetalIncludesHost.h | 61 +++ .../metal/GPUReconstructionMetalKernels.mm | 80 ++++ ...PUReconstructionMetalKernelsSpecialize.inc | 27 ++ GPU/GPUTracking/CMakeLists.txt | 68 ++- dependencies/FindO2GPU.cmake | 53 ++- 9 files changed, 911 insertions(+), 36 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..9c9b7e4fd2b15 --- /dev/null +++ b/GPU/GPUTracking/Base/metal/CMakeLists.txt @@ -0,0 +1,110 @@ +# 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 is the first version with a generic address space: earlier versions +# reject an unannotated pointer or `this` outright, which GPUCommonDefAPI.h +# relies on for GPUgeneric() and GPUdDefault(). +set(METAL_FLAGS -std=metal4.1 ${GPUCA_METAL_DENORMALS_FLAGS}) +if(GPUCA_DETERMINISTIC_MODE GREATER_EQUAL ${GPUCA_DETERMINISTIC_MODE_MAP_NO_FAST_MATH}) + set(METAL_FLAGS ${METAL_FLAGS} ${GPUCA_METAL_NO_FAST_MATH_FLAGS}) +endif() +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 +) + +set(SRCS GPUReconstructionMetal.mm GPUReconstructionMetalKernels.mm) +set(HDRS GPUReconstructionMetal.h GPUReconstructionMetalIncludesHost.h) + +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 + + # 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 compiled library into __DATA,__gpu_resource during final link. This + # way we do not need to create an intermediate object. Compiling the source at + # run time is not an option: the driver's compiler service dies on it. + target_link_options(${targetName} + PRIVATE + "-Wl,-sectcreate,__DATA,__gpu_resource,${METAL_BIN}.ir") + 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..47e64045a596a --- /dev/null +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal @@ -0,0 +1,77 @@ +// 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 +// MSL rejects derived classes outside of this pragma, and the kernels are +// class-based throughout. metal_stdlib itself uses it in 46 paired places. +#pragma METAL internals : enable +using namespace metal; + +// --- OpenCL compatibility shims --------------------------------------------- + +// Address space aliases (match OpenCL vernacular used by the project); constant +// is spelled the same in MSL +#define global device +#define local threadgroup + +#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) + +// --- double ------------------------------------------------------------------ +// MSL has no double. GPUdoubleBinary64 is IEEE-754 binary64 in software, with the +// same eight bytes in the same order, so the keyword can simply name it and the +// shared code needs no separate spelling. Must come after metal_stdlib, which +// uses the token itself. +#include "GPUCommonDoubleBinary64.h" +#define double o2::gpu::GPUdoubleBinary64 +#include "GPUCommonDouble.h" + +// --- Project headers --------------------------------------------------------- +#include "GPUCommonDef.h" +#include "GPUCommonTypeTraits.h" +#include "GPUCommonArray.h" + +#include "GPUConstantMem.h" +#include "GPUReconstructionIncludesDeviceAll.h" + +// --- Kernel list expansion --------------------------------------------------- +#define GPUCA_KRNL(...) GPUCA_KRNLGPU(__VA_ARGS__) + +// --- Constant memory + global heap plumbing --------------------------------- +// The heap and the constant memory arrive as buffer(0) and buffer(1). The latter +// is untyped because a buffer of GPUConstantMem, which has base classes, is not +// a valid kernel argument type. +#define GPUCA_CONSMEM_PTR \ + device char* gpu_mem [[buffer(0)]], \ + device char* pConstantRaw [[buffer(1)]], +#define GPUCA_CONSMEM (*(device GPUConstantMem*)pConstantRaw) + +#include "GPUReconstructionKernelList.h" + +// 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..53a114976cf68 --- /dev/null +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.h @@ -0,0 +1,66 @@ +// 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 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; + + 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..18fa84dfbe688 --- /dev/null +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm @@ -0,0 +1,405 @@ +// 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() +{ + 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"); + } + + // FIXME: until I understand how to enable the gGPUConstantMemBufferSize constexpr + int32_t tmpGPUConstantMemBufferSize = 100000; // gGPUConstantMemBufferSize + mInternals->mem_constant = [mInternals->device newBufferWithLength:tmpGPUConstantMemBufferSize options:MTLResourceStorageModeShared]; + if (mInternals->mem_constant == nil) { + 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() +{ + // Tear down Metal + 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; + + 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 %p", src); + } + + if (dst >= mHostMemoryBase && dst < ((char*)mHostMemoryBase + mHostMemorySize)) { + destBuffer = mInternals->mem_host; + destOffset = (char*)dst - (char*)mHostMemoryBase; + } else if (dst >= mDeviceMemoryBase && dst < ((char*)mDeviceMemoryBase + mDeviceMemorySize)) { + destBuffer = mInternals->mem_gpu; + destOffset = (char*)dst - (char*)mDeviceMemoryBase; + } else { + GPUErrorReturn("Unknown buffer at %p", dst); + } + + [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 %p", 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]; +} + +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) +{ + 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() +{ + GPUInfo("Loading Metal library (Platform version %s)", [mInternals->device.architecture.name cStringUsingEncoding:NSUTF8StringEncoding]); + + unsigned long sz = 0; + const uint8_t* p = find_gpu_resource_in_image("libO2GPUTrackingMETAL.dylib", &sz); + if (p == nullptr || sz == 0) { + GPUError("Metal library not found in the __gpu_resource section"); + return 1; + } + + // the section is part of the mapped image, so it outlives the dispatch_data_t + // and does not have to be copied + dispatch_data_t blob = dispatch_data_create(p, sz, nullptr, ^{ + }); + + NSError* error = nil; + mInternals->library = [mInternals->device newLibraryWithData:blob error:&error]; + + if (mInternals->library == nil) { + NSLog(@"%@", error); + GPUError("Error loading the Metal library"); + 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..8a2c6eacd1545 --- /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, constants::GPU_MAX_STREAMS> commandQueues; + std::array, constants::GPU_MAX_STREAMS> commandBuffers; + + std::vector> functions; + std::vector> pipelines; + + id mem_gpu; + id mem_constant; + id mem_host; + + id library; +}; +} // 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..7fe46f1395aac --- /dev/null +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm @@ -0,0 +1,80 @@ +// 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; + + 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..cc44b2003a01d 100644 --- a/GPU/GPUTracking/CMakeLists.txt +++ b/GPU/GPUTracking/CMakeLists.txt @@ -463,41 +463,39 @@ if (onnxruntime_FOUND) 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() +if(METAL_ENABLED) + add_subdirectory(Base/metal) +elseif(CUDA_ENABLED OR OPENCL_ENABLED OR HIP_ENABLED) + 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 955f333d8e412..cf342e00c68be 100644 --- a/dependencies/FindO2GPU.cmake +++ b/dependencies/FindO2GPU.cmake @@ -10,7 +10,7 @@ # or submit itself to any jurisdiction. # NOTE!!!! - Whenever this file is changed, move it over to alidist/resources -# FindO2GPU.cmake Version 20 +# FindO2GPU.cmake Version 21 set(CUDA_COMPUTETARGET_DEFAULT_FULL 80-real;86-real;89-real;120-real;75-virtual) set(HIP_AMDGPUTARGET_DEFAULT_FULL gfx906;gfx908) @@ -20,6 +20,10 @@ set(HIP_AMDGPUTARGET_DEFAULT_MINIMAL gfx906) if(NOT DEFINED ENABLE_CUDA) set(ENABLE_CUDA "AUTO") endif() +if(NOT DEFINED ENABLE_METAL) + # OFF rather than AUTO while the backend is unproven + set(ENABLE_METAL "OFF") +endif() if(NOT DEFINED ENABLE_OPENCL) set(ENABLE_OPENCL "AUTO") endif() @@ -28,7 +32,16 @@ if(NOT DEFINED ENABLE_HIP) endif() string(TOUPPER "${ENABLE_CUDA}" ENABLE_CUDA) string(TOUPPER "${ENABLE_OPENCL}" ENABLE_OPENCL) +string(TOUPPER "${ENABLE_METAL}" ENABLE_METAL) string(TOUPPER "${ENABLE_HIP}" ENABLE_HIP) +if(APPLE) + # macOS ships OpenCL 1.2 only, below the 2.x the OpenCL backend needs; Metal + # replaces it there. + if(ENABLE_OPENCL AND NOT ENABLE_OPENCL STREQUAL "AUTO") + message(FATAL_ERROR "The OpenCL backend needs OpenCL 2.x and macOS ships 1.2; the Metal backend replaces it there") + endif() + set(ENABLE_OPENCL "OFF") +endif() if(NOT DEFINED CMAKE_BUILD_TYPE_UPPER) string(TOUPPER "${CMAKE_BUILD_TYPE}" CMAKE_BUILD_TYPE_UPPER) endif() @@ -149,6 +162,7 @@ if(GPUCA_DETERMINISTIC_NO_FTZ) set(GPUCA_CXX_DENORMALS_FLAGS "") set(GPUCA_CUDA_DENORMALS_FLAGS "--ftz=false") set(GPUCA_OCL_DENORMALS_FLAGS "") + set(GPUCA_METAL_DENORMALS_FLAGS "") set(GPUCA_HIP_DENORMALS_FLAGS "-fno-gpu-flush-denormals-to-zero") else() if (CMAKE_SYSTEM_NAME MATCHES Darwin OR NOT CMAKE_SYSTEM_PROCESSOR MATCHES "(x86)|(X86)|(amd64)|(AMD64)") @@ -158,11 +172,13 @@ else() endif() set(GPUCA_CUDA_DENORMALS_FLAGS "--ftz=true") set(GPUCA_OCL_DENORMALS_FLAGS "-cl-denorms-are-zero") + set(GPUCA_METAL_DENORMALS_FLAGS "") set(GPUCA_HIP_DENORMALS_FLAGS "-fgpu-flush-denormals-to-zero") endif() set(GPUCA_CXX_NO_FAST_MATH_FLAGS "-fno-fast-math -ffp-contract=off") set(GPUCA_CUDA_NO_FAST_MATH_FLAGS "--prec-div=true --prec-sqrt=true --fmad false -Xcompiler -fno-fast-math -Xcompiler -ffp-contract=off") set(GPUCA_OCL_NO_FAST_MATH_FLAGS -cl-fp32-correctly-rounded-divide-sqrt ) +set(GPUCA_METAL_NO_FAST_MATH_FLAGS "-fno-fast-math -ffp-contract=off") if(GPUCA_DETERMINISTIC_MODE GREATER_EQUAL ${GPUCA_DETERMINISTIC_MODE_MAP_WHOLEO2}) add_definitions(-DGPUCA_DETERMINISTIC_MODE) string(APPEND CMAKE_CXX_FLAGS_${CMAKE_BUILD_TYPE_UPPER} " ${GPUCA_CXX_NO_FAST_MATH_FLAGS}") @@ -429,6 +445,41 @@ if(ENABLE_HIP) endif() endif() +# =================================== Metal ================================== +if(ENABLE_METAL) + find_library(METAL_FRAMEWORK Metal) + find_library(COREFOUNDATION_FRAMEWORK CoreFoundation) + find_library(FOUNDATION_FRAMEWORK Foundation) + find_library(QUARTZCORE_FRAMEWORK QuartzCore) + # The frameworks are there on every macOS, but the backend needs MSL 4.1, which + # is macOS 27 and its toolchain, so ask the compiler instead of assuming. A + # library built as MSL 4.1 also only loads on macOS 27 and later. + if(NOT DEFINED GPUCA_METAL_MSL41) + set(GPUCA_METAL_PROBE "${CMAKE_CURRENT_BINARY_DIR}/metal_msl41_probe.metal") + file(WRITE "${GPUCA_METAL_PROBE}" "#include \nkernel void probe(device float* o [[buffer(0)]], uint i [[thread_position_in_grid]]) { o[i] = o[i] * 2.0f; }\n") + execute_process(COMMAND xcrun -sdk macosx metal -std=metal4.1 -c "${GPUCA_METAL_PROBE}" -o "${GPUCA_METAL_PROBE}.air" + RESULT_VARIABLE GPUCA_METAL_PROBE_RESULT OUTPUT_QUIET ERROR_QUIET) + if(GPUCA_METAL_PROBE_RESULT EQUAL 0) + set(GPUCA_METAL_MSL41 ON CACHE INTERNAL "Metal toolchain compiles MSL 4.1") + else() + set(GPUCA_METAL_MSL41 OFF CACHE INTERNAL "Metal toolchain compiles MSL 4.1") + endif() + endif() + if(METAL_FRAMEWORK AND COREFOUNDATION_FRAMEWORK AND FOUNDATION_FRAMEWORK AND QUARTZCORE_FRAMEWORK AND GPUCA_METAL_MSL41) + set(METAL_ENABLED ON) + set(METAL_FRAMEWORKS ${METAL_FRAMEWORK} ${COREFOUNDATION_FRAMEWORK} + ${FOUNDATION_FRAMEWORK} ${QUARTZCORE_FRAMEWORK}) + message(STATUS "Found Metal frameworks") + elseif(NOT ENABLE_METAL STREQUAL "AUTO") + if(NOT METAL_FRAMEWORK OR NOT COREFOUNDATION_FRAMEWORK OR NOT FOUNDATION_FRAMEWORK OR NOT QUARTZCORE_FRAMEWORK) + message(FATAL_ERROR "Metal frameworks not available") + endif() + message(FATAL_ERROR "The Metal backend needs MSL 4.1: this toolchain rejects -std=metal4.1, and the result would need macOS 27 to run") + else() + set(METAL_ENABLED OFF) + endif() +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)