GPU Metal: wait for and record the events of kernels and copies

The kernels, the memory clear, GPUMemCpy and WriteToConstantMemory
ignored the events they were given, which is how the chain orders the
work of its lanes across streams. Wait for the event list before and
record the event after, as the CUDA backend does.

Author: ktf

GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm

@@ -232,6 +232,9 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const
   }
 
   auto realStream = stream == -1 ? 0 : stream;
+  if (nEvents) {
+    StreamWaitForEvents(realStream, evList, nEvents);
+  }
   const MetalMemRef s = metalResolve(mInternals, src, mHostMemoryBase, mHostMemorySize, mDeviceMemoryBase, mDeviceMemorySize);
   const MetalMemRef d = metalResolve(mInternals, dst, mHostMemoryBase, mHostMemorySize, mDeviceMemoryBase, mDeviceMemorySize);
 
@@ -252,6 +255,9 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const
     [blit endEncoding];
     CommitStream(realStream);
   }
+  if (ev) {
+    RecordMarker(ev, realStream);
+  }
 
   if (GetProcessingSettings().serializeGPU & 2) {
     GPUDebug(("GPUMemCpy " + std::to_string(toGPU)).c_str(), stream, true);
@@ -282,6 +288,9 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const
     [blit endEncoding];
     CommitStream(realStream);
   }
+  if (ev) {
+    RecordMarker(ev, realStream);
+  }
 
   if (GetProcessingSettings().serializeGPU & 2) {
     GPUDebug("WriteToConstantMemory", stream, true);

GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm

@@ -25,6 +25,10 @@
   const uint32_t num = GetKernelNum<T, I>();
   auto& kExec = _xyz.x;
   auto& runRange = _xyz.y;
+  auto& kEvents = _xyz.z;
+  if (kEvents.evList) {
+    StreamWaitForEvents(kExec.stream, kEvents.evList, kEvents.nEvents);
+  }
 
   id<MTLComputePipelineState> pso = mInternals->pipelines[num];
   id<MTLComputeCommandEncoder> computeEncoder = [mInternals->commandBuffers[kExec.stream] computeCommandEncoder];
@@ -51,6 +55,9 @@
   // without committing the kernel never reaches the device.
   [computeEncoder endEncoding];
   CommitStream(kExec.stream);
+  if (kEvents.ev) {
+    RecordMarker(kEvents.ev, kExec.stream);
+  }
 }
 
 template <class T, int32_t I>

GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernelsSpecialize.inc

@@ -16,6 +16,9 @@ inline void GPUReconstructionMetal::runKernelBackend<GPUMemClean16, 0>(const krn
 {
   const uint64_t offset = static_cast<const uint8_t*>(ptr) - static_cast<const uint8_t*>(mDeviceMemoryBase);
   const uint64_t length = (size + 15ull) & ~15ull;
+  if (_xyz.z.evList) {
+    StreamWaitForEvents(_xyz.x.stream, _xyz.z.evList, _xyz.z.nEvents);
+  }
 
   id<MTLCommandBuffer> cb = mInternals->commandBuffers[_xyz.x.stream];
   id<MTLBlitCommandEncoder> blit = [cb blitCommandEncoder];
@@ -24,4 +27,7 @@ inline void GPUReconstructionMetal::runKernelBackend<GPUMemClean16, 0>(const krn
   [blit endEncoding];
 
   CommitStream(_xyz.x.stream);
+  if (_xyz.z.ev) {
+    RecordMarker(_xyz.z.ev, _xyz.x.stream);
+  }
 }