From 32841da271faa6842d8503af5e68a40d9257230c Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Tue, 10 Nov 2026 18:25:05 +0100 Subject: [PATCH 1/3] GPU Metal: wait for synchronous copies to complete A copy or constant-memory write on stream -1 is synchronous in the CUDA and HIP backends, and the callers read the destination right after it. The Metal backend only encoded the copy into stream 0's open command buffer, so e.g. CheckErrorCodes read a stale error buffer and lost the errors raised after the last synchronisation, such as ERROR_MERGER_HIT_OVERFLOW. Drain the stream before returning. --- GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm index 5292eb5b2f15e..093a4b83b87b7 100644 --- a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm @@ -336,6 +336,9 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const if (ev) { RecordMarker(ev, realStream); } + if (stream == -1) { // Synchronous like cudaMemcpy: the caller reads the destination right away + DrainStream(realStream); + } if (GetProcessingSettings().serializeGPU & 2) { GPUDebug(("GPUMemCpy " + std::to_string(toGPU)).c_str(), stream, true); @@ -370,6 +373,9 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const if (ev) { RecordMarker(ev, realStream); } + if (stream == -1) { + DrainStream(realStream); + } if (GetProcessingSettings().serializeGPU & 2) { GPUDebug("WriteToConstantMemory", stream, true); From a31a726e91c36eb8c346d9b4877cdc35c564ed98 Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Tue, 10 Nov 2026 18:25:06 +0100 Subject: [PATCH 2/3] GPU Metal: keep the memory pools resident for the kernels The kernels reach the device and host pools through the addresses stored in constant memory, so Metal does not know a dispatch uses them and does not keep them mapped for the GPU. About a second after the last blit that touched a pool, the kernels read it as empty: every loop bounded by a count in device memory finishes in ~20 us, Finalize step 2 takes 13 ms instead of 0.4 ms, and the next blit stalls 0.2-0.3 s before it starts, mapping the pool again. At 20 kHz the merger TrackFit alone runs for a second, so everything queued behind it did nothing until the next transfer: 0 looper candidates, 1,246 of 21,679 retry refits, 536 of 2,701 looper fits, 245,145 tracks for 233,521 on the CPU at 27 sectors. It is not the length of the dispatch: at pbpb30 a 1.5 s host sleep before a normal TrackFit gives the same picture (0 retry refits, 0 looper candidates, 22,250 tracks for 21,360), and a 1.7 s single TrackFit dispatch completes once the pools are made resident with useResource (207 looper fits, 4,521 retry refits, 10,775 candidates, 21,361 tracks, as without the sleep). pbpb30: 985-1087 ms/event against 1002-1075 before (3 rounds of 10 events); the debug dumps are byte-identical, sim5pp identical to the CPU. --- GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm index 09780ac104232..f8c1fbeb2a743 100644 --- a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm @@ -42,6 +42,10 @@ // Slots 0 to 2 are what GPUCA_KRNL_BUFFERS and GPUCA_KRNL_SECTOR_ARG declare // in the generated signature; the kernel's own arguments follow from slot 3. [computeEncoder setComputePipelineState:pso]; + // The kernels reach the pools through the addresses kept in constant memory, + // which Metal cannot see, so they have to be made resident explicitly. + [computeEncoder useResource:mInternals->mem_gpu usage:MTLResourceUsageRead | MTLResourceUsageWrite]; + [computeEncoder useResource:mInternals->mem_host usage:MTLResourceUsageRead | MTLResourceUsageWrite]; [computeEncoder setBuffer:mInternals->mem_gpu offset:0 atIndex:0]; [computeEncoder setBuffer:mInternals->mem_constant offset:0 atIndex:1]; const int32_t iSector = runRange.index; From 8a18d96473314ad5dca9624ee127293d8a0ee69d Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Tue, 10 Nov 2026 18:25:10 +0100 Subject: [PATCH 3/3] GPU Metal: optionally fail when a kernel ran on an unmapped memory pool A pool the GPU cannot see is read as zeros, so the kernels return without an error and the chain goes on with results they did not produce. With GPUCA_METAL_CANARY=1, a one-thread kernel after each dispatch reads, through the pool's address like the kernels do, a word the host wrote past the end of the device pool, and counts a read that does not return it in a buffer bound directly; the drain then fails the run, as it does for a failed command buffer. Without the previous commit, the pbpb30 reproducer (1.5 s host sleep before TrackFit) ends with "14 kernels ran while the memory pool was not mapped for the GPU" and a GPU Backend Failure instead of 22,250 tracks. It is opt-in because the extra dispatch costs 1021-1038 ms/event against 979-989 at pbpb30 (3 rounds of 10 events). --- .../Base/metal/GPUReconstructionMETAL.metal | 9 ++++++++ .../Base/metal/GPUReconstructionMetal.mm | 22 ++++++++++++++++++- .../GPUReconstructionMetalIncludesHost.h | 16 ++++++++++++++ .../metal/GPUReconstructionMetalKernels.mm | 3 +++ 4 files changed, 49 insertions(+), 1 deletion(-) diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal b/GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal index 14f7986f606a5..5d6dc49b30e5f 100644 --- a/GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal @@ -82,5 +82,14 @@ using namespace metal; #include "GPUReconstructionKernelList.h" +// Reads a word of the memory pool through its address, as the kernels do, and +// counts a read that does not return it; the host fails the run on a count. +kernel void krnl_metalCanary(device atomic_uint* bad [[buffer(0)]], constant uint64_t& word [[buffer(1)]]) +{ + if (*(device uint32_t*)word != 0xA11CE0A5u) { + atomic_fetch_add_explicit(bad, 1u, memory_order_relaxed); + } +} + // clang-format on #pragma clang diagnostic pop diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm index 093a4b83b87b7..2f8a28da23922 100644 --- a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm @@ -185,7 +185,12 @@ uint32_t metalGPUCoreCount() // The kernels reconstruct device pointers from these bases, so they must be // the addresses the GPU sees, not the CPU-visible ones of a shared buffer. mDeviceMemoryBase = (void*)(uintptr_t)mInternals->mem_gpu.gpuAddress; - mDeviceMemorySize = mInternals->mem_gpu.allocatedSize; + mDeviceMemorySize = mInternals->mem_gpu.allocatedSize - sizeof(uint32_t); + *(uint32_t*)((char*)mInternals->mem_gpu.contents + mDeviceMemorySize) = mInternals->canaryMagic; + if (getenv("GPUCA_METAL_CANARY")) { + mInternals->canaryBad = [mInternals->device newBufferWithLength:sizeof(uint32_t) options:MTLResourceStorageModeShared]; + *(uint32_t*)mInternals->canaryBad.contents = 0; + } mDeviceConstantMem = (GPUConstantMem*)(uintptr_t)mInternals->mem_constant.gpuAddress; if (GetProcessingSettings().debugLevel >= 1) { @@ -233,6 +238,7 @@ uint32_t metalGPUCoreCount() if (mDeviceMemoryBase) { [mInternals->mem_gpu release]; [mInternals->mem_constant release]; + [mInternals->canaryPipeline release]; for (uint32_t i = 0; i < mInternals->functions.size(); i++) { [mInternals->functions[i] release]; } @@ -252,6 +258,7 @@ uint32_t metalGPUCoreCount() [mInternals->streamFence[i] release]; } [mInternals->mem_host release]; + [mInternals->canaryBad release]; } // The completion handlers of the limit still use the internals std::unique_lock lock(mInternals->inFlightMutex); @@ -476,6 +483,10 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const [cb release]; } mInternals->submittedBuffers[stream].clear(); + if (const uint32_t bad = mInternals->canaryBad ? *(volatile uint32_t*)mInternals->canaryBad.contents : 0) { + GPUError("Metal: %u kernels ran while the memory pool was not mapped for the GPU", bad); + error = error ? error : 1; + } if (drainStart) { fprintf(mInternals->timeline, "DR s%d %.6f %.6f\n", stream, drainStart, metalHostTime()); } @@ -612,5 +623,14 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const return 1; } + id canary = [mInternals->library newFunctionWithName:@"krnl_metalCanary"]; + mInternals->canaryPipeline = canary ? [mInternals->device newComputePipelineStateWithFunction:canary error:&error] : nil; + [canary release]; + if (mInternals->canaryPipeline == nil) { + NSLog(@"%@", error); + GPUError("Error creating the Metal canary pipeline state"); + return 1; + } + return AddKernels(); } diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h index e17275bdbbc2a..23b2e851de76f 100644 --- a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h @@ -96,6 +96,12 @@ struct GPUReconstructionMetalInternals { id mem_gpu; id mem_constant; id mem_host; + // Opt-in, GPUCA_METAL_CANARY=1: a kernel that reads a word past the end of the + // device pool, where the host put canaryMagic, and counts the reads that return + // something else. canaryBad is nil when it is off. + static constexpr uint32_t canaryMagic = 0xA11CE0A5; + id canaryPipeline; + id canaryBad = nil; id library; @@ -156,6 +162,16 @@ inline void metalTimelineAdd(GPUReconstructionMetalInternals* internals, int32_t internals->timelineOps[stream] += buf; } } +// Runs after each kernel in its encoder, so a kernel that ran while the pool +// was not mapped for the GPU is counted and the run fails instead of going on +// with the results it did not produce. +inline void metalEncodeCanary(GPUReconstructionMetalInternals* internals, id encoder, uint64_t word) +{ + [encoder setComputePipelineState:internals->canaryPipeline]; + [encoder setBuffer:internals->canaryBad offset:0 atIndex:0]; + [encoder setBytes:&word length:sizeof(word) atIndex:1]; + [encoder dispatchThreadgroups:MTLSizeMake(1, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)]; +} template inline void metalEndEncoder(GPUReconstructionMetalInternals* internals, E encoder, int32_t stream) { diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm index f8c1fbeb2a743..b9c532485cba1 100644 --- a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm @@ -65,6 +65,9 @@ } [computeEncoder dispatchThreadgroups:MTLSizeMake(kExec.nBlocks, 1, 1) threadsPerThreadgroup:MTLSizeMake(threadGroupSize, 1, 1)]; + if (mInternals->canaryBad) { + metalEncodeCanary(mInternals, computeEncoder, (uint64_t)mDeviceMemoryBase + mDeviceMemorySize); + } metalEndEncoder(mInternals, computeEncoder, kExec.stream); id cb = [mInternals->commandBuffers[kExec.stream] retain]; if (kEvents.ev) {