GPU CMake: enable the Metal backend by default where the toolchain supports it

GPU/GPUTracking/Base/GPUReconstructionTimeframe.cxx

@@ -290,8 +290,8 @@ int32_t GPUReconstructionTimeframe::LoadCreateTimeFrame(int32_t iEvent)
             return (1);
           }
           nTotalClusters += nClusters;
-          printf("Placing event %4d+%d (ID %4d) at z %7.3f (time %'dns) %s(collisions %4d, bunch %6ld, train %3d) (%'10d clusters, %'10d MC labels, %'10d track MC info)\n", nCollisions, nBorderCollisions, useEvent, shift, (int32_t)(nBunch * config.bunchSpacing), inTF ? " inside" : "outside",
-                 nCollisions, nBunch, nTrain, nClusters, mChain->mIOPtrs.nMCLabelsTPC, mChain->mIOPtrs.nMCInfosTPC);
+          printf("Placing event %4d+%d (ID %4d) at z %7.3f (time %'dns) %s(collisions %4d, bunch %6lld, train %3d) (%'10d clusters, %'10d MC labels, %'10d track MC info)\n", nCollisions, nBorderCollisions, useEvent, shift, (int32_t)(nBunch * config.bunchSpacing), inTF ? " inside" : "outside",
+                 nCollisions, (long long)nBunch, nTrain, nClusters, mChain->mIOPtrs.nMCLabelsTPC, mChain->mIOPtrs.nMCInfosTPC);
           nInBunchPileUp++;
           nCollisionsInTrain++;
           p2 *= mCollisionProbability / nInBunchPileUp;

GPU/GPUTracking/CMakeLists.txt

@@ -513,9 +513,7 @@ if(CPUDEVICE_ENABLED)
 endif()
 
 if(ALIGPU_BUILD_TYPE STREQUAL "O2" OR ALIGPU_BUILD_TYPE STREQUAL "Standalone")
-  if(NOT ${CMAKE_SYSTEM} MATCHES Darwin)
-    add_subdirectory(Standalone/Benchmark)
-  endif()
+  add_subdirectory(Standalone/Benchmark)
 endif()
 
 if(GPUCA_DETERMINISTIC_MODE GREATER_EQUAL ${GPUCA_DETERMINISTIC_MODE_MAP_GPU})

GPU/GPUTracking/Definitions/GPUSettingsList.h

@@ -348,8 +348,8 @@ AddOption(keepAllMemory, bool, false, "", 0, "Allocate all memory on both device
 AddOption(keepDisplayMemory, bool, false, "", 0, "Like keepAllMemory, but only for memory required for event display")
 AddOption(disableMemoryReuse, bool, false, "", 0, "Disable memory reusage (for debugging only)")
 AddOption(memoryAllocationStrategy, int8_t, 0, "", 0, "Memory Allocation Strategy (0 = auto, 1 = individual allocations, 2 = single global allocation)")
-AddOption(forceMemoryPoolSize, uint64_t, 1, "memSize", 0, "Force size of allocated GPU / page locked host memory (0 = disable, 1 / 2 = auto, other values set the size of the memory pool)", min(0ul))
-AddOption(forceHostMemoryPoolSize, uint64_t, 0, "hostMemSize", 0, "Force size of allocated host page locked host memory (overriding memSize)", min(0ul))
+AddOption(forceMemoryPoolSize, uint64_t, 1, "memSize", 0, "Force size of allocated GPU / page locked host memory (0 = disable, 1 / 2 = auto, other values set the size of the memory pool)", min(uint64_t(0)))
+AddOption(forceHostMemoryPoolSize, uint64_t, 0, "hostMemSize", 0, "Force size of allocated host page locked host memory (overriding memSize)", min(uint64_t(0)))
 AddOption(memoryScalingFactor, float, 1.f, "", 0, "Factor to apply to all memory scalers")
 AddOption(memoryScalingFuzz, uint64_t, 0, "", 0, "Fuzz the memoryScalingFactor (0 disable, 1 enable, >1 set seed", def(1))
 AddOption(tpcInputWithClusterRejection, uint8_t, 0, "", 0, "Indicate whether the TPC input is CTF data with cluster rejection, to tune buffer estimations")

GPU/GPUTracking/Definitions/GPUSettingsListStandalone.h

@@ -47,7 +47,7 @@ EndConfig()
 BeginConfig(GPUSettingsStandalone, configStandalone)
 AddOption(runGPU, uint8_t, 1, "", 'g', "Use GPU for processing", message("GPU processing enabled"), set(2))
 AddOptionSet(runGPU, uint8_t, 0, "", 'c', "Use CPU for processing", message("CPU enabled"))
-AddOption(gpuType, std::string, "AUTO", "", 0, "GPU type (CUDA / HIP / OCL / CPUDEVICE) or CPU or AUTO")
+AddOption(gpuType, std::string, "AUTO", "", 0, "GPU type (CUDA / HIP / OCL / METAL / CPUDEVICE) or CPU or AUTO")
 AddOption(runGPUforce, bool, true, "", 0, "Force usage of the specified GPU device type, no CPU fallback")
 AddOption(promptOnExit, bool, false, "", 0, "Prompt for keypress before exiting")
 AddOption(continueOnError, bool, false, "", 0, "Continue processing after an error")
@@ -64,8 +64,8 @@ AddOption(eventDisplay, int32_t, 0, "display", 'd', "Show standalone event displ
 AddOption(eventGenerator, bool, false, "", 0, "Run event generator")
 AddOption(cont, bool, false, "", 0, "Process continuous timeframe data, even if input is triggered")
 AddOption(setMaxTimeBin, int32_t, -2, "", 0, "maximum time bin of continuous data, 0 for triggered events, -1 for automatic continuous mode, -2 for automatic continuous / triggered")
-AddOption(outputcontrolmem, uint64_t, 0, "outputMemory", 0, "Use predefined output buffer of this size", min(0ul), message("Using %s bytes as output memory"))
-AddOption(inputcontrolmem, uint64_t, 0, "inputMemory", 0, "Use predefined input buffer of this size", min(0ul), message("Using %s bytes as input memory"))
+AddOption(outputcontrolmem, uint64_t, 0, "outputMemory", 0, "Use predefined output buffer of this size", min(uint64_t(0)), message("Using %s bytes as output memory"))
+AddOption(inputcontrolmem, uint64_t, 0, "inputMemory", 0, "Use predefined input buffer of this size", min(uint64_t(0)), message("Using %s bytes as input memory"))
 AddOption(cpuAffinity, int32_t, -1, "", 0, "Pin CPU affinity to this CPU core", min(-1))
 AddOption(fifoScheduler, bool, false, "", 0, "Use FIFO realtime scheduler", message("Setting FIFO scheduler: %s"))
 AddOption(fpe, int8_t, -1, "", 0, "Trap on floating point exceptions (-1 = if no ffast-math)")

GPU/GPUTracking/Standalone/Benchmark/standalone.cxx

@@ -124,6 +124,7 @@ int32_t ReadConfiguration(int argc, char** argv)
   setlocale(LC_ALL, "en_US.utf-8");
   setlocale(LC_NUMERIC, "en_US.utf-8");
   if (configStandalone.cpuAffinity != -1) {
+#ifdef __linux__
     cpu_set_t mask;
     CPU_ZERO(&mask);
     CPU_SET(configStandalone.cpuAffinity, &mask);
@@ -133,8 +134,13 @@ int32_t ReadConfiguration(int argc, char** argv)
       printf("Error setting CPU affinity\n");
       return 1;
     }
+#else
+    printf("CPU affinity not supported on this platform\n");
+    return 1;
+#endif
   }
   if (configStandalone.fifoScheduler) {
+#ifdef __linux__
     printf("Setting FIFO scheduler\n");
     sched_param param;
     sched_getparam(0, &param);
@@ -143,13 +149,21 @@ int32_t ReadConfiguration(int argc, char** argv)
       printf("Error setting scheduler\n");
       return 1;
     }
+#else
+    printf("FIFO scheduler not supported on this platform\n");
+    return 1;
+#endif
   }
 #ifdef __FAST_MATH__
   if (configStandalone.fpe == 1) {
 #else
   if (configStandalone.fpe) {
 #endif
+#ifdef __linux__
     feenableexcept(FE_INVALID | FE_DIVBYZERO | FE_OVERFLOW);
+#else
+    printf("Trapping floating point exceptions not supported on this platform\n");
+#endif
   }
   bool detMode = false, noFTZMode = false;
 #ifdef GPUCA_DETERMINISTIC_MODE
@@ -1010,7 +1024,7 @@ int32_t main(int argc, char** argv)
       }
     }
     if (nEventsProcessed > 1) {
-      printf("Total: %ld clusters, %ld tracks\n", nClustersTotal, nTracksTotal);
+      printf("Total: %lld clusters, %lld tracks\n", (long long)nClustersTotal, (long long)nTracksTotal);
     }
   }
 
@@ -1021,7 +1035,9 @@ int32_t main(int argc, char** argv)
 
 #ifndef _WIN32
   if (configStandalone.proc.runQA && configStandalone.fpe) {
+#ifdef __linux__
     fedisableexcept(FE_INVALID | FE_DIVBYZERO | FE_OVERFLOW);
+#endif
   }
 #endif
 

dependencies/FindO2GPU.cmake

@@ -10,7 +10,7 @@
 # or submit itself to any jurisdiction.
 
 # NOTE!!!! - Whenever this file is changed, move it over to alidist/resources
-# FindO2GPU.cmake Version 23
+# FindO2GPU.cmake Version 24
 
 set(CUDA_COMPUTETARGET_DEFAULT_FULL 80-real;86-real;89-real;120-real;75-virtual)
 set(HIP_AMDGPUTARGET_DEFAULT_FULL gfx906;gfx908)
@@ -21,8 +21,7 @@ if(NOT DEFINED ENABLE_CUDA)
   set(ENABLE_CUDA "AUTO")
 endif()
 if(NOT DEFINED ENABLE_METAL)
-  # OFF rather than AUTO while the backend is unproven
-  set(ENABLE_METAL "OFF")
+  set(ENABLE_METAL "AUTO")
 endif()
 if(NOT DEFINED ENABLE_OPENCL)
   set(ENABLE_OPENCL "AUTO")
@@ -467,15 +466,14 @@ if(ENABLE_METAL)
   # The frameworks are there on every macOS, but the backend needs MSL 4.1, which
   # is macOS 27 and its toolchain, so ask the compiler instead of assuming. A
   # library built as MSL 4.1 also only loads on macOS 27 and later.
-  if(NOT DEFINED GPUCA_METAL_MSL41)
+  # Only a success is cached, so that installing the toolchain is picked up.
+  if(NOT GPUCA_METAL_MSL41)
     set(GPUCA_METAL_PROBE "${CMAKE_CURRENT_BINARY_DIR}/metal_msl41_probe.metal")
     file(WRITE "${GPUCA_METAL_PROBE}" "#include <metal_stdlib>\nkernel void probe(device float* o [[buffer(0)]], uint i [[thread_position_in_grid]]) { o[i] = o[i] * 2.0f; }\n")
     execute_process(COMMAND xcrun -sdk macosx metal -std=metal4.1 -c "${GPUCA_METAL_PROBE}" -o "${GPUCA_METAL_PROBE}.air"
-                    RESULT_VARIABLE GPUCA_METAL_PROBE_RESULT OUTPUT_QUIET ERROR_QUIET)
+                    RESULT_VARIABLE GPUCA_METAL_PROBE_RESULT OUTPUT_QUIET ERROR_VARIABLE GPUCA_METAL_PROBE_ERROR)
     if(GPUCA_METAL_PROBE_RESULT EQUAL 0)
       set(GPUCA_METAL_MSL41 ON CACHE INTERNAL "Metal toolchain compiles MSL 4.1")
-    else()
-      set(GPUCA_METAL_MSL41 OFF CACHE INTERNAL "Metal toolchain compiles MSL 4.1")
     endif()
   endif()
   if(METAL_FRAMEWORK AND COREFOUNDATION_FRAMEWORK AND FOUNDATION_FRAMEWORK AND QUARTZCORE_FRAMEWORK AND GPUCA_METAL_MSL41)
@@ -487,8 +485,9 @@ if(ENABLE_METAL)
     if(NOT METAL_FRAMEWORK OR NOT COREFOUNDATION_FRAMEWORK OR NOT FOUNDATION_FRAMEWORK OR NOT QUARTZCORE_FRAMEWORK)
       message(FATAL_ERROR "Metal frameworks not available")
     endif()
-    message(FATAL_ERROR "The Metal backend needs MSL 4.1: this toolchain rejects -std=metal4.1, and the result would need macOS 27 to run")
+    message(FATAL_ERROR "The Metal backend needs MSL 4.1: this toolchain rejects -std=metal4.1, and the result would need macOS 27 to run\n${GPUCA_METAL_PROBE_ERROR}")
   else()
+    message(STATUS "Metal backend disabled, the toolchain does not compile MSL 4.1: ${GPUCA_METAL_PROBE_ERROR}")
     set(METAL_ENABLED OFF)
   endif()
 endif()

prodtests/full_system_test.sh

@@ -48,6 +48,7 @@ NJOBS=${NJOBS:-"${NCPUS}"}
 SHMSIZE=${SHMSIZE:-8000000000} # Size of shared memory for messages (use 128 GB for 550 event full TF)
 #TPCTRACKERSCRATCHMEMORY=${TPCTRACKERSCRATCHMEMORY:-8000000000} # Not needed by default
 ENABLE_GPU_TEST=${ENABLE_GPU_TEST:-0} # Run the full system test also on the GPU
+FST_GPUTYPE=${FST_GPUTYPE:-CUDA} # GPU backend used in case ENABLE_GPU_TEST=1 (CUDA / HIP / METAL)
 GPUMEMSIZE=${GPUMEMSIZE:-6000000000} # Size of GPU memory to use in case ENABBLE_GPU_TEST=1
 NTIMEFRAMES=${NTIMEFRAMES:-1} # Number of time frames to process
 TFDELAY=${TFDELAY:-100} # Delay in seconds between publishing time frames
@@ -281,7 +282,7 @@ for STAGE in $STAGES; do
   export WORKFLOW_PARAMETERS=
   if [[ "$STAGE" = "WITHGPU" ]]; then
     export CREATECTFDICT=0
-    export GPUTYPE=CUDA
+    export GPUTYPE=$FST_GPUTYPE
     export HOSTMEMSIZE=1000000000
     export SYNCMODE=1
     export CTFINPUT=0