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)