From 49e884f3e8934ee4c99d84e22f0d7ed0164da068 Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Thu, 5 Nov 2026 15:50:35 +0100 Subject: [PATCH 1/5] GPU Metal: bound the threadgroup size when creating pipelines Pass the tuned thread count as maxTotalThreadsPerThreadgroup, the Metal equivalent of __launch_bounds__. Without it the register allocation can leave room for fewer threads than the tuning table assumes, and the launch was clamped: on an M1 Max (applegpu_g13s) the two TRD tracker kernels got 384 instead of 512, and the TPC compression gather kernel 832 instead of 1024, which breaks its block-wide scans. --- .../Base/metal/GPUReconstructionMetalKernels.mm | 15 +++++++++++---- 1 file changed, 11 insertions(+), 4 deletions(-) diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm index af3897b727c7c..e4a858220588b 100644 --- a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm @@ -64,17 +64,24 @@ return 1; } + // The kernel sizes its threadgroup memory from the tuning table at compile + // time, so a launch must not exceed what it was built for either way. + const uint32_t tunedThreads = this->template getKernelProperties().nThreads; + + // Like __launch_bounds__: without it the register allocation may leave room + // for fewer threads than the kernel is tuned for. + MTLComputePipelineDescriptor* desc = [MTLComputePipelineDescriptor new]; + desc.computeFunction = krnl; + desc.maxTotalThreadsPerThreadgroup = tunedThreads; NSError* error = nil; - id pso = [mInternals->device newComputePipelineStateWithFunction:krnl error:&error]; + id pso = [mInternals->device newComputePipelineStateWithDescriptor:desc options:MTLPipelineOptionNone reflection:nil error:&error]; + [desc release]; if (pso == nil) { NSLog(@"%@", error); GPUError("Error creating Metal pipeline state: %s", [kname cStringUsingEncoding:NSUTF8StringEncoding]); return 1; } - // The kernel sizes its threadgroup memory from the tuning table at compile - // time, so a launch must not exceed what it was built for either way. - const uint32_t tunedThreads = this->template getKernelProperties().nThreads; if (tunedThreads > pso.maxTotalThreadsPerThreadgroup) { // Fewer threads than the kernel was built for: safe for the shared arrays, // but the scan and compaction kernels assume every slot has an owner. From 25286d82fede21ef2036947c68d4e888e903e0c3 Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Thu, 5 Nov 2026 14:22:31 +0100 Subject: [PATCH 2/5] GPU Metal: commit the memory clear through CommitStream The GPUMemClean16 blit committed its command buffer directly, so the stream kept the committed buffer as its open one, and the next encoder on that stream aborted with "failed assertion _status < MTLCommandBufferStatusCommitted". The clear was also never waited for. --- .../Base/metal/GPUReconstructionMetalKernelsSpecialize.inc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernelsSpecialize.inc b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernelsSpecialize.inc index 1ee192cff822f..7c2abe7cf8954 100644 --- a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernelsSpecialize.inc +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernelsSpecialize.inc @@ -23,5 +23,5 @@ inline void GPUReconstructionMetal::runKernelBackend(const krn [blit fillBuffer:mInternals->mem_gpu range:NSMakeRange(offset, length) value:0]; [blit endEncoding]; - [cb commit]; + CommitStream(_xyz.x.stream); } From d8f88b90648ef69fe690d8d0f31ab8325682461a Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Thu, 5 Nov 2026 14:22:31 +0100 Subject: [PATCH 3/5] GPU Metal: create the device events The event slots were only placement-constructed as nil ids, so the first StreamWaitForEvents threw "object cannot be nil" from encodeWaitForEvent. Create an MTLSharedEvent per slot, as the CUDA backend creates its events, and release them on exit. --- .../Base/metal/GPUReconstructionMetal.mm | 14 +++++++++++++- 1 file changed, 13 insertions(+), 1 deletion(-) diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm index 821a314d7d83a..0aef2048e26a7 100644 --- a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm @@ -137,7 +137,12 @@ for (uint32_t i = 0; i < mEvents.size(); i++) { auto* events = (id*)mEvents[i].data(); - new (events) id[ mEvents[i].size() ]; + for (uint32_t j = 0; j < mEvents[i].size(); j++) { + events[j] = [mInternals->device newSharedEvent]; + if (events[j] == nil) { + GPUErrorReturn("Error creating Metal event"); + } + } } return (0); @@ -148,6 +153,13 @@ // Tear down Metal SynchronizeGPU(); + for (uint32_t i = 0; i < mEvents.size(); i++) { + auto* events = (id*)mEvents[i].data(); + for (uint32_t j = 0; j < mEvents[i].size(); j++) { + [events[j] release]; + } + } + if (mMaster == nullptr) { if (mDeviceMemoryBase) { [mInternals->mem_gpu release]; From e31ce9048abf72e14918f502740ea749c5a798f5 Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Thu, 5 Nov 2026 14:45:39 +0100 Subject: [PATCH 4/5] GPU Metal: give the device events CUDA's semantics The events signalled and waited for the fixed value 1, and ReleaseEvent reset them to 0, so a wait issued after a release, or for an event not recorded yet, never completed: the GPU watchdog killed each such command buffer after about 10 s, and the queue then dropped all later work. On a 5-event pp time frame this made the TPC reconstruction take 80 s and return nothing. Count up per event instead: RecordMarker signals the next value, a wait asks for the last recorded one and is a no-op if there is none, as cudaStreamWaitEvent is, and ReleaseEvent does nothing, as in CUDA. --- .../Base/metal/GPUReconstructionMetal.h | 1 + .../Base/metal/GPUReconstructionMetal.mm | 20 ++++++++++++------- .../GPUReconstructionMetalIncludesHost.h | 8 ++++++++ 3 files changed, 22 insertions(+), 7 deletions(-) diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.h b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.h index df571fa3c1525..50eb40ebf7ccc 100644 --- a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.h +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.h @@ -59,6 +59,7 @@ class GPUReconstructionMetal : public GPUReconstructionProcessing::KernelInterfa void CommitStream(int32_t stream); void DrainStream(int32_t stream); + uint64_t EventValue(deviceEvent ev, bool next); private: int32_t AddKernels(); diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm index 0aef2048e26a7..6a305b5e73fb2 100644 --- a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm @@ -291,10 +291,14 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const 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]; + // Nothing to do, as in CUDA: the next RecordMarker supersedes the last one. +} + +uint64_t GPUReconstructionMetal::EventValue(deviceEvent ev, bool next) +{ + std::lock_guard lock(mInternals->eventMutex); + auto& value = mInternals->eventValues[ev.get()]; + return next ? ++value : value; } void GPUReconstructionMetal::RecordMarker(deviceEvent* ev, int32_t stream) @@ -303,7 +307,7 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const // 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 encodeSignalEvent:mtlEvent value:EventValue(*ev, true)]; CommitStream(stream); } @@ -353,7 +357,9 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const 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]; + if (const uint64_t value = EventValue(evList[ei], false)) { + [cb encodeWaitForEvent:mtlEvent value:value]; + } } CommitStream(stream); DrainStream(stream); @@ -363,7 +369,7 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const { for (int32_t i = 0; i < nEvents; i++) { auto mtlEvent = (__bridge id)(evList[i].get()); - if (mtlEvent.signaledValue == 0) { + if (mtlEvent.signaledValue < EventValue(evList[i], false)) { return false; } } diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h index 2d93f5b08adfb..9c32ba0dc4bae 100644 --- a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h @@ -32,6 +32,8 @@ using half = _Float16; #include #include #include +#include +#include #include "GPULogging.h" #include "GPUReconstructionMetal.h" @@ -67,6 +69,12 @@ struct GPUReconstructionMetalInternals { id mem_host; id library; + + // The value each event was last recorded with. An MTLSharedEvent only counts + // up, so recording signals the next value and waiting asks for the last one, + // which makes waiting for a never recorded event a no-op, as in CUDA. + std::unordered_map eventValues; + std::mutex eventMutex; }; } // namespace o2::gpu From 5aa49f0605f96fd73dc396b2cb7231eb83e83904 Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Thu, 5 Nov 2026 14:45:40 +0100 Subject: [PATCH 5/5] GPU Metal: report command buffers that failed on the device DrainStream waited for the submitted command buffers but never looked at their status, so a kernel that faulted or hung only showed up as missing output. Print the Metal error of a failed buffer. --- GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm | 3 +++ 1 file changed, 3 insertions(+) diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm index 6a305b5e73fb2..00cd5d8687624 100644 --- a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm @@ -327,6 +327,9 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const // committed, so it never completes. Only what was submitted can be waited for. for (id cb : mInternals->submittedBuffers[stream]) { [cb waitUntilCompleted]; + if (cb.status == MTLCommandBufferStatusError) { + GPUError("Metal command buffer on stream %d failed: %s", stream, [[cb.error description] UTF8String]); + } [cb release]; } mInternals->submittedBuffers[stream].clear();