GPU Metal: implement the device kernel timers

With --debug 1 and deviceTimers (the default), the kernel timers expect the
backend to report the kernel's GPU time, and Metal reported none, so all
kernel times read 0. Give the timed kernel a command buffer of its own and
report its GPUEndTime - GPUStartTime. The times are only set once the
completion is processed, so wait for it rather than for the stream's event.

pbpb30 on M1 Max: 1012 ms of kernel time per event, vs 1218 ms on the CPU.

Author: ktf

GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm

@@ -29,6 +29,11 @@
   if (kEvents.evList) {
     StreamWaitForEvents(kExec.stream, kEvents.evList, kEvents.nEvents);
   }
+  // Device timers: the kernel gets a command buffer of its own, whose GPU time is the kernel's
+  const bool timed = GetProcessingSettings().deviceTimers && GetProcessingSettings().debugLevel > 0;
+  if (timed && mInternals->streamHasWork[kExec.stream]) {
+    CommitStream(kExec.stream);
+  }
 
   id<MTLComputePipelineState> pso = mInternals->pipelines[num];
   id<MTLComputeCommandEncoder> computeEncoder = [mInternals->commandBuffers[kExec.stream] computeCommandEncoder];
@@ -57,9 +62,19 @@
   [computeEncoder dispatchThreadgroups:MTLSizeMake(kExec.nBlocks, 1, 1)
                  threadsPerThreadgroup:MTLSizeMake(threadGroupSize, 1, 1)];
   metalEndEncoder(mInternals, computeEncoder, kExec.stream);
+  id<MTLCommandBuffer> cb = [mInternals->commandBuffers[kExec.stream] retain];
   if (kEvents.ev) {
     RecordMarker(kEvents.ev, kExec.stream);
   }
+  if (timed) {
+    if (mInternals->commandBuffers[kExec.stream] == cb) {
+      CommitStream(kExec.stream);
+    }
+    DrainStream(kExec.stream);
+    [cb waitUntilCompleted];  // the times are only set once the completion is processed
+    _xyz.t = cb.GPUEndTime - cb.GPUStartTime;
+  }
+  [cb release];
 }
 
 template <class T, int32_t I>