GPU Metal: reconstruction on Metal bringup

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();

GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm

@@ -137,7 +137,12 @@
 
   for (uint32_t i = 0; i < mEvents.size(); i++) {
     auto* events = (id<MTLSharedEvent>*)mEvents[i].data();
-    new (events) id<MTLSharedEvent>[ 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<MTLSharedEvent>*)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];
@@ -279,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<MTLSharedEvent>)(ev.get<void*>());
-  [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<std::mutex> lock(mInternals->eventMutex);
+  auto& value = mInternals->eventValues[ev.get<void*>()];
+  return next ? ++value : value;
 }
 
 void GPUReconstructionMetal::RecordMarker(deviceEvent* ev, int32_t stream)
@@ -291,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<MTLSharedEvent>)(ev->get<void*>());
-  [cb encodeSignalEvent:mtlEvent value:1];
+  [cb encodeSignalEvent:mtlEvent value:EventValue(*ev, true)];
   CommitStream(stream);
 }
 
@@ -311,6 +327,9 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const
   // committed, so it never completes. Only what was submitted can be waited for.
   for (id<MTLCommandBuffer> 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();
@@ -341,7 +360,9 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const
   id<MTLCommandBuffer> cb = mInternals->commandBuffers[stream];
   for (int32_t ei = 0; ei < nEvents; ei++) {
     auto mtlEvent = (__bridge id<MTLSharedEvent>)(evList[ei].get<void*>());
-    [cb encodeWaitForEvent:mtlEvent value:1];
+    if (const uint64_t value = EventValue(evList[ei], false)) {
+      [cb encodeWaitForEvent:mtlEvent value:value];
+    }
   }
   CommitStream(stream);
   DrainStream(stream);
@@ -351,7 +372,7 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const
 {
   for (int32_t i = 0; i < nEvents; i++) {
     auto mtlEvent = (__bridge id<MTLSharedEvent>)(evList[i].get<void*>());
-    if (mtlEvent.signaledValue == 0) {
+    if (mtlEvent.signaledValue < EventValue(evList[i], false)) {
       return false;
     }
   }

GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h

@@ -32,6 +32,8 @@ using half = _Float16;
 #include <vector>
 #include <string>
 #include <memory>
+#include <mutex>
+#include <unordered_map>
 #include "GPULogging.h"
 
 #include "GPUReconstructionMetal.h"
@@ -67,6 +69,12 @@ struct GPUReconstructionMetalInternals {
   id<MTLBuffer> mem_host;
 
   id<MTLLibrary> 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<void*, uint64_t> eventValues;
+  std::mutex eventMutex;
 };
 } // namespace o2::gpu
 

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<T, I>().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<MTLComputePipelineState> pso = [mInternals->device newComputePipelineStateWithFunction:krnl error:&error];
+  id<MTLComputePipelineState> 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<T, I>().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.

GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernelsSpecialize.inc

@@ -23,5 +23,5 @@ inline void GPUReconstructionMetal::runKernelBackend<GPUMemClean16, 0>(const krn
   [blit fillBuffer:mInternals->mem_gpu range:NSMakeRange(offset, length) value:0];
   [blit endEncoding];
 
-  [cb commit];
+  CommitStream(_xyz.x.stream);
 }