GPU: smaller thread groups for the merger track fit and the tracking data on Metal

GPU/GPUTracking/Base/metal/CMakeLists.txt

@@ -28,7 +28,9 @@ set(METAL_BIN ${CMAKE_CURRENT_BINARY_DIR}/GPUReconstructionMetalCode)
 # relies on for GPUgeneric() and GPUdDefault().
 set(METAL_FLAGS -std=metal4.1 ${GPUCA_METAL_DENORMALS_FLAGS})
 if(GPUCA_DETERMINISTIC_MODE GREATER_EQUAL ${GPUCA_DETERMINISTIC_MODE_MAP_NO_FAST_MATH})
-  set(METAL_FLAGS ${METAL_FLAGS} ${GPUCA_METAL_NO_FAST_MATH_FLAGS})
+  # A string of several flags, which the custom commands below would pass as one
+  separate_arguments(GPUCA_METAL_NO_FAST_MATH_FLAGS_LIST UNIX_COMMAND "${GPUCA_METAL_NO_FAST_MATH_FLAGS}")
+  set(METAL_FLAGS ${METAL_FLAGS} ${GPUCA_METAL_NO_FAST_MATH_FLAGS_LIST})
 endif()
 set(METAL_DEFINES "-D$<JOIN:$<TARGET_PROPERTY:O2::GPUTracking,COMPILE_DEFINITIONS>,$<SEMICOLON>-D>"
             "-I$<JOIN:$<FILTER:$<TARGET_PROPERTY:O2::GPUTracking,INCLUDE_DIRECTORIES>,EXCLUDE,^/usr/include/?>,$<SEMICOLON>-I>"
@@ -46,7 +48,7 @@ if(ALIGPU_BUILD_TYPE STREQUAL "O2")
                  PUBLIC_LINK_LIBRARIES O2::GPUTracking
                  TARGETVARNAME targetName)
 
-  target_link_libraries(${targetName} PUBLIC ${METAL_FRAMEWORKS})
+  target_link_libraries(${targetName} PUBLIC ${METAL_FRAMEWORKS} "-framework IOKit")
 
   target_compile_definitions(${targetName} PRIVATE $<TARGET_PROPERTY:O2::GPUTracking,COMPILE_DEFINITIONS>)
   # the compile_defitions are not propagated automatically on purpose (they are

GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm

@@ -18,6 +18,7 @@
 #include <cstddef>
 #include <map>
 
+#include <IOKit/IOKitLib.h>
 #include <mach-o/dyld.h>
 #include <mach-o/getsect.h>
 #include <mach-o/ldsyms.h>  // _mh_execute_header
@@ -28,6 +29,28 @@
     return (1);             \
   }
 
+namespace
+{
+// Metal does not report the number of GPU cores, but the IORegistry entry of
+// an Apple silicon GPU does. 0 if it is not there.
+uint32_t metalGPUCoreCount()
+{
+  uint32_t n = 0;
+  io_service_t service = IOServiceGetMatchingService(kIOMainPortDefault, IOServiceMatching("AGXAccelerator"));
+  if (service) {
+    CFTypeRef p = IORegistryEntrySearchCFProperty(service, kIOServicePlane, CFSTR("gpu-core-count"), kCFAllocatorDefault, kIORegistryIterateRecursively);
+    if (p) {
+      if (CFGetTypeID(p) == CFNumberGetTypeID()) {
+        CFNumberGetValue((CFNumberRef)p, kCFNumberSInt32Type, &n);
+      }
+      CFRelease(p);
+    }
+    IOObjectRelease(service);
+  }
+  return n;
+}
+}  // namespace
+
 #include "utils/qGetLdBinarySymbols.h"
 QGET_LD_BINARY_SYMBOLS(GPUReconstructionMetalCode_src);
 
@@ -74,6 +97,14 @@
     mDeviceName = device_name.c_str();
     // Basically a random number for now.
     mMaxBackendThreads = 1000;
+    // The automatic grids are minBlocks per core; with the default of 1 they
+    // left all but one core idle.
+    mMultiprocessorCount = metalGPUCoreCount();
+    if (mMultiprocessorCount == 0) {
+      GPUWarning("Could not read the number of GPU cores, assuming 1");
+      mMultiprocessorCount = 1;
+    }
+    mMaxBackendThreads = std::max<int32_t>(mMaxBackendThreads, deviceMaxWorkGroup.width * mMultiprocessorCount);
 
     if (GetMetalPrograms()) {
       return 1;
@@ -105,6 +136,10 @@
       if (mInternals->commandBuffers[i] == nil) {
         GPUErrorReturn("Error creating Metal command buffer");
       }
+      mInternals->streamOrder[i] = [mInternals->device newSharedEvent];
+      mInternals->streamOrderValue[i] = 0;
+      mInternals->streamFence[i] = [mInternals->device newFence];
+      mInternals->streamHasWork[i] = false;
     }
 
     mInternals->mem_host = [mInternals->device newBufferWithLength:mHostMemorySize options:MTLResourceStorageModeShared];
@@ -179,6 +214,8 @@
         DrainStream(i);
         [mInternals->commandBuffers[i] release];
         [mInternals->commandQueues[i] release];
+        [mInternals->streamOrder[i] release];
+        [mInternals->streamFence[i] release];
       }
       [mInternals->mem_host release];
     }
@@ -232,6 +269,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);
 
@@ -244,13 +284,16 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const
   } else {
     id<MTLCommandBuffer> cb = mInternals->commandBuffers[realStream];
     id<MTLBlitCommandEncoder> blit = [cb blitCommandEncoder];
+    metalBeginEncoder(mInternals, blit, realStream);
     [blit copyFromBuffer:s.buffer
             sourceOffset:s.offset
                 toBuffer:d.buffer
        destinationOffset:d.offset
                     size:sizeBytes];
-    [blit endEncoding];
-    CommitStream(realStream);
+    metalEndEncoder(mInternals, blit, realStream);
+  }
+  if (ev) {
+    RecordMarker(ev, realStream);
   }
 
   if (GetProcessingSettings().serializeGPU & 2) {
@@ -274,13 +317,16 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const
   } else {
     id<MTLCommandBuffer> cb = mInternals->commandBuffers[realStream];
     id<MTLBlitCommandEncoder> blit = [cb blitCommandEncoder];
+    metalBeginEncoder(mInternals, blit, realStream);
     [blit copyFromBuffer:s.buffer
             sourceOffset:s.offset
                 toBuffer:mInternals->mem_constant
        destinationOffset:offset
                     size:size];
-    [blit endEncoding];
-    CommitStream(realStream);
+    metalEndEncoder(mInternals, blit, realStream);
+  }
+  if (ev) {
+    RecordMarker(ev, realStream);
   }
 
   if (GetProcessingSettings().serializeGPU & 2) {
@@ -314,15 +360,25 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const
 void GPUReconstructionMetal::CommitStream(int32_t stream)
 {
   id<MTLCommandBuffer> cb = mInternals->commandBuffers[stream];
+  // A CUDA stream runs its work in order, a Metal queue only as far as its
+  // hazard tracking sees, and the kernels reach most memory through addresses
+  // in constant memory. So each buffer of a stream waits for the previous one.
+  const uint64_t value = ++mInternals->streamOrderValue[stream];
+  [cb encodeSignalEvent:mInternals->streamOrder[stream] value:value];
   [cb commit];
   // Hand the committed buffer over to the stream's submission list, which owns
   // it until it is waited for, and open a fresh one to encode the next work in.
   mInternals->submittedBuffers[stream].emplace_back(cb);
+  mInternals->streamHasWork[stream] = false;
   mInternals->commandBuffers[stream] = [[mInternals->commandQueues[stream] commandBuffer] retain];
+  [mInternals->commandBuffers[stream] encodeWaitForEvent:mInternals->streamOrder[stream] value:value];
 }
 
 void GPUReconstructionMetal::DrainStream(int32_t stream)
 {
+  if (mInternals->streamHasWork[stream]) {
+    CommitStream(stream);
+  }
   // Waiting on the buffer that is currently open would block forever: it is not
   // committed, so it never completes. Only what was submitted can be waited for.
   for (id<MTLCommandBuffer> cb : mInternals->submittedBuffers[stream]) {

GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h

@@ -59,6 +59,13 @@ struct GPUReconstructionMetalInternals {
   // encoded into afterwards, so the two cannot be the same object.
   std::array<id<MTLCommandBuffer>, constants::GPU_MAX_STREAMS> commandBuffers;
   std::array<std::vector<id<MTLCommandBuffer>>, constants::GPU_MAX_STREAMS> submittedBuffers;
+  // Signalled by each committed buffer and waited for by the next one, see CommitStream
+  std::array<id<MTLSharedEvent>, constants::GPU_MAX_STREAMS> streamOrder;
+  std::array<uint64_t, constants::GPU_MAX_STREAMS> streamOrderValue;
+  // Orders the encoders within the open command buffer of a stream, which
+  // collects the stream's work until something has to wait for it.
+  std::array<id<MTLFence>, constants::GPU_MAX_STREAMS> streamFence;
+  std::array<bool, constants::GPU_MAX_STREAMS> streamHasWork;
 
   std::vector<id<MTLFunction>> functions;
   std::vector<id<MTLComputePipelineState>> pipelines;
@@ -76,6 +83,20 @@ struct GPUReconstructionMetalInternals {
   std::unordered_map<void*, uint64_t> eventValues;
   std::mutex eventMutex;
 };
+// An encoder of a stream starts after the previous one of the same stream and
+// leaves its work in the open command buffer.
+template <class E>
+inline void metalBeginEncoder(GPUReconstructionMetalInternals* internals, E encoder, int32_t stream)
+{
+  [encoder waitForFence:internals->streamFence[stream]];
+}
+template <class E>
+inline void metalEndEncoder(GPUReconstructionMetalInternals* internals, E encoder, int32_t stream)
+{
+  [encoder updateFence:internals->streamFence[stream]];
+  [encoder endEncoding];
+  internals->streamHasWork[stream] = true;
+}
 } // namespace o2::gpu
 
 #endif // GPURECONSTRUCTIONOMETALINCLUDESHOST_H

GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm

@@ -25,9 +25,14 @@
   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];
+  metalBeginEncoder(mInternals, computeEncoder, kExec.stream);
 
   // 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.
@@ -47,10 +52,10 @@
   }
   [computeEncoder dispatchThreadgroups:MTLSizeMake(kExec.nBlocks, 1, 1)
                  threadsPerThreadgroup:MTLSizeMake(threadGroupSize, 1, 1)];
-  // Without ending the encoder the next one on this buffer is rejected, and
-  // without committing the kernel never reaches the device.
-  [computeEncoder endEncoding];
-  CommitStream(kExec.stream);
+  metalEndEncoder(mInternals, computeEncoder, kExec.stream);
+  if (kEvents.ev) {
+    RecordMarker(kEvents.ev, kExec.stream);
+  }
 }
 
 template <class T, int32_t I>

GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernelsSpecialize.inc

@@ -16,12 +16,17 @@ 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];
+  metalBeginEncoder(mInternals, blit, _xyz.x.stream);
 
   [blit fillBuffer:mInternals->mem_gpu range:NSMakeRange(offset, length) value:0];
-  [blit endEncoding];
-
-  CommitStream(_xyz.x.stream);
+  metalEndEncoder(mInternals, blit, _xyz.x.stream);
+  if (_xyz.z.ev) {
+    RecordMarker(_xyz.z.ev, _xyz.x.stream);
+  }
 }

GPU/GPUTracking/Definitions/Parameters/GPUParameters.csv

@@ -5,7 +5,7 @@ WARP_SIZE,0,,64,64,32,32,32,32,32,32,32,32,32,32,32,64,32,64,32
 THREAD_COUNT_DEFAULT,256,,256,256,,,,,,512,512,,512,256,512,512,512,,256
 ,,,,,,,,,,,,,,,,,,,
 LB:,,,,,,,,,,,,,,,,,,,
-GPUTPCCreateTrackingData,256,,"[256, 7]","[192, 2]",,,,,,"[224, 7]",256,"[128, 14]",416,,"[64, 21]",,384,"[320, 2]",
+GPUTPCCreateTrackingData,256,,"[256, 7]","[192, 2]",,,,,,"[224, 7]",256,"[128, 14]",416,,"[64, 21]",,384,"[320, 2]",64
 GPUTPCTrackletConstructor,256,,"[768, 8]","[512, 10]","[256, 2]","[256, 1]","[256, 2]","[1024, 2]","[512, 4]",1024,"[256, 2]",1024,"[1024, 1]",,"[768, 2]",,768,512,
 GPUTPCTrackletSelector,256,,"[384, 5]","[192, 10]","[256, 3]","[256, 1]","[256, 3]","[512, 4]","[256, 3]","[288, 3]","[192, 3]","[544, 1]","[32, 2]",,"[384, 3]",,992,"[256, 6]",
 GPUTPCNeighboursFinder,256,,"[192, 8]","[960, 8]",256,256,256,512,256,864,"[640, 1]","[512, 2]","[736, 1]",,"[480, 3]",,992,"[704, 1]",
@@ -16,7 +16,7 @@ GPUTPCCreateOccupancyMap_fill,256,,,,,,,,,256,,256,256,,256,,,256,
 GPUTPCCreateOccupancyMap_fold,256,,,,,,,,,256,,256,256,,256,,,256,
 GPUTRDTrackerKernels_o2Version,512,,,,,,,,,512,,512,512,,512,,,512,
 GPUTPCCompressionKernels_step0attached,256,,"[128, 1]","[64, 2]",,,,,,"[160, 2]",128,"[448, 1]",352,,"[1024, 1]",,"[96, 3]","[128, 4]",
-GPUTPCCompressionKernels_step1unattached,256,,"[512, 2]","[512, 2]",,,,,,"[288, 4]","[512, 2]","[256, 4]","[512, 2]",,"[512, 3]",,"[512, 2]","[512, 3]",
+GPUTPCCompressionKernels_step1unattached,256,,"[512, 2]","[512, 2]",,,,,,"[288, 4]","[512, 2]","[256, 4]","[512, 2]",,"[512, 3]",,"[512, 2]","[512, 3]","[256, 8]"
 GPUTPCDecompressionKernels_step0attached,256,,"[128, 2]","[128, 2]",,,,,,"[32, 1]","[32, 1]","[32, 1]","[32, 1]",,"[128, 1]",,"[32, 1]","[128, 1]",
 GPUTPCDecompressionKernels_step1unattached,256,,"[64, 2]","[64, 2]",,,,,,"[32, 1]","[32, 1]","[32, 1]","[32, 1]",,"[64, 1]",,"[32, 1]","[64, 1]",
 GPUTPCDecompressionUtilKernels_sortPerSectorRow,256,,,,,,,,,256,,256,256,,256,,,256,
@@ -27,7 +27,7 @@ GPUTPCCFDecodeZSLink,"""GPUCA_WARP_SIZE""",,"""GPUCA_WARP_SIZE""","""GPUCA_WARP_
 GPUTPCCFDecodeZSDenseLink,"""GPUCA_WARP_SIZE""",,"[""GPUCA_WARP_SIZE"", 4]","[""GPUCA_WARP_SIZE"", 14]",,,,,,"[""GPUCA_WARP_SIZE"", 14]","""GPUCA_WARP_SIZE""","[""GPUCA_WARP_SIZE"", 22]","[""GPUCA_WARP_SIZE"", 22]",,"[""GPUCA_WARP_SIZE"", 17]",,"[""GPUCA_WARP_SIZE"", 8]","[""GPUCA_WARP_SIZE"", 5]",
 GPUTPCCFGather,"[1024, 1]",,"[1024, 5]","[1024, 1]",,,,,,"[160, 11]","[1024, 1]",736,896,,"[928, 1]",,"[1024, 1]","[320, 2]",
 COMPRESSION_GATHER,1024,,1024,1024,,,,,,1024,1024,,1024,,,,,,
-GPUTPCGMMergerTrackFit,256,,"[192, 2]","[64, 7]",,,,,,"[32, 16]","[32, 8]","[32, 14]","[160, 2]",,"[32, 24]",,"[64, 8]","[64, 6]",
+GPUTPCGMMergerTrackFit,256,,"[192, 2]","[64, 7]",,,,,,"[32, 16]","[32, 8]","[32, 14]","[160, 2]",,"[32, 24]",,"[64, 8]","[64, 6]","[64, 4]"
 GPUTPCGMMergerFollowLoopers,256,,"[256, 5]","[256, 4]",,,,,,"[256, 4]","[128, 4]","[1024, 1]",640,,"[128, 16]",,"[224, 3]","[256, 7]",
 GPUTPCGMMergerSectorRefit,256,,"[64, 4]","[256, 2]",,,,,,"[32, 8]","[64, 5]","[32, 7]","[32, 7]",,"[32, 20]",,"[32, 10]","[64, 4]",
 GPUTPCGMMergerUnpackResetIds,256,,256,256,,,,,,256,256,256,256,,256,,256,256,