davidrohr · Oct 11, 2026
@@ -17,8 +17,10 @@
 
 #include "GPUCommonDef.h"
 
-#if !defined(GPUCA_GPUCODE) // Could also enable custom search on the CPU, but it is not always faster, so we stick to std::sort
+#if !defined(GPUCA_GPUCODE) || defined(__CPUDEVICE__) // Could also enable custom search on the CPU, but it is not always faster, so we stick to std::sort
+#ifndef GPUCA_GPUCODE_DEVICE
 #include <algorithm>
+#endif
 #define GPUCA_ALGORITHM_STD
 #endif
 
@@ -280,7 +282,7 @@ GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end GPUCA_THREAD_INFO_
 template <class T, class S>
 GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end, const S& comp GPUCA_THREAD_INFO_DECL)
 {
-#ifndef GPUCA_GPUCODE
+#if !defined(GPUCA_GPUCODE) || defined(GPUCA_DEVICE_SINGLE_THREAD)
   GPUCommonAlgorithm::sort(begin, end, comp);
 #elif defined(GPUCA_DETERMINISTIC_MODE) // Not using GPUCA_DETERMINISTIC_CODE, which is enforced in TPC compression
   if (get_local_id(0) == 0) {

@@ -15,7 +15,7 @@
 #ifndef GPUCOMMONARRAY_H
 #define GPUCOMMONARRAY_H
 
-#if !defined(GPUCA_GPUCODE_DEVICE) || defined(__CUDACC__) || defined(__HIPCC__) // TODO: Get rid of GPUCommonArray once OpenCL supports <array>
+#if !defined(GPUCA_GPUCODE_DEVICE) || defined(__CUDACC__) || defined(__HIPCC__) || defined(__CPUDEVICE__) // TODO: Get rid of GPUCommonArray once OpenCL supports <array>
 #ifndef GPUCA_GPUCODE_COMPILEKERNELS
 #include <array>
 #endif

@@ -31,12 +31,12 @@
 #include "GPUCommonDefSettings.h"
 
 #if !defined(__CLING__) && !defined(G__ROOT) // No GPU code for ROOT
-  #if defined(__CUDACC__) || defined(__OPENCL__) || defined(__HIPCC__) || defined(__OPENCL_HOST__) || defined(__METAL__) || defined(__METAL_HOST__)
+  #if defined(__CUDACC__) || defined(__OPENCL__) || defined(__HIPCC__) || defined(__OPENCL_HOST__) || defined(__METAL__) || defined(__METAL_HOST__) || defined(__CPUDEVICE__)
     #define GPUCA_GPUCODE // Compiled by GPU compiler
   #endif
 
   #if defined(GPUCA_GPUCODE)
-    #if defined(__CUDA_ARCH__) || defined(__OPENCL__) || defined(__HIP_DEVICE_COMPILE__) || defined(__METAL_VERSION__)
+    #if defined(__CUDA_ARCH__) || defined(__OPENCL__) || defined(__HIP_DEVICE_COMPILE__) || defined(__METAL_VERSION__) || defined(__CPUDEVICE__)
       #define GPUCA_GPUCODE_DEVICE // Executed on device
     #endif
     #if defined(__CUDACC__)
@@ -47,6 +47,8 @@
       #define GPUCA_GPUTYPE OCL
     #elif defined(__METAL__) || defined(__METAL_HOST__)
       #define GPUCA_GPUTYPE METAL
+    #elif defined(__CPUDEVICE__)
+      #define GPUCA_GPUTYPE CPUDEVICE
     #endif
   #endif
 #endif

@@ -27,7 +27,7 @@
 //Define macros for GPU keywords. i-version defines inline functions.
 //All host-functions in GPU code are automatically inlined, to avoid duplicate symbols.
 //For non-inline host only functions, use no keyword at all!
-#if !defined(GPUCA_GPUCODE) || defined(__OPENCL_HOST__) || defined(__METAL_HOST__) // For host / ROOT dictionary
+#if !defined(GPUCA_GPUCODE) || defined(__OPENCL_HOST__) || defined(__METAL_HOST__) || defined(__CPUDEVICE__) // For host / ROOT dictionary, and for the CPUDevice backend whose device code runs on the host
   #define GPUd()                                    // device function
   #define GPUdDefault()                             // default (constructor / operator) device function
   #define GPUhdDefault()                            // default (constructor / operator) host device function

@@ -236,7 +236,7 @@ GPUdi() constexpr T GPUCommonMath::nextMultipleOf(T val)
 
 GPUdi() float2 GPUCommonMath::MakeFloat2(float x, float y)
 {
-#if !defined(GPUCA_GPUCODE) || defined(__OPENCL__) || defined(__OPENCL_HOST__) || defined(__METAL__) || defined(__METAL_HOST__)
+#if !defined(GPUCA_GPUCODE) || defined(__OPENCL__) || defined(__OPENCL_HOST__) || defined(__METAL__) || defined(__METAL_HOST__) || defined(__CPUDEVICE__)
   float2 ret = {x, y};
   return ret;
 #else

@@ -427,7 +427,7 @@ class GPUReconstruction
     void* mGPULib;
     void* mGPUEntry;
   };
-  static std::shared_ptr<LibraryLoader> sLibCUDA, sLibHIP, sLibOCL, sLibMETAL;
+  static std::shared_ptr<LibraryLoader> sLibCUDA, sLibHIP, sLibOCL, sLibMETAL, sLibCPUDEVICE;
 
   // Debugging
   struct debugInternal;

@@ -17,5 +17,6 @@
 #cmakedefine HIP_ENABLED
 #cmakedefine OPENCL_ENABLED
 #cmakedefine METAL_ENABLED
+#cmakedefine CPUDEVICE_ENABLED
 #cmakedefine GPUCA_COMPILER_VERSIONS @GPUCA_COMPILER_VERSIONS@
 // clang-format on
@@ -106,6 +106,10 @@ std::shared_ptr<GPUReconstruction::LibraryLoader>* GPUReconstruction::GetLibrary
   } else if (type == DeviceType::METAL) {
 #ifdef METAL_ENABLED
     return &sLibMETAL;
+#endif
+  } else if (type == DeviceType::CPUDEVICE) {
+#ifdef CPUDEVICE_ENABLED
+    return &sLibCPUDEVICE;
 #endif
   } else {
     GPUError("Error: Invalid device type %u", (uint32_t)type);
@@ -131,6 +135,7 @@ std::shared_ptr<GPUReconstruction::LibraryLoader> GPUReconstruction::sLibCUDA(ne
 std::shared_ptr<GPUReconstruction::LibraryLoader> GPUReconstruction::sLibHIP(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingHIP" LIBRARY_EXTENSION, "GPUReconstruction_Create_HIP"));
 std::shared_ptr<GPUReconstruction::LibraryLoader> GPUReconstruction::sLibOCL(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingOCL" LIBRARY_EXTENSION, "GPUReconstruction_Create_OCL"));
 std::shared_ptr<GPUReconstruction::LibraryLoader> GPUReconstruction::sLibMETAL(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingMETAL" LIBRARY_EXTENSION, "GPUReconstruction_Create_METAL"));
+std::shared_ptr<GPUReconstruction::LibraryLoader> GPUReconstruction::sLibCPUDEVICE(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingCPUDevice" LIBRARY_EXTENSION, "GPUReconstruction_Create_CPUDEVICE"));
 
 GPUReconstruction::LibraryLoader::LibraryLoader(const char* lib, const char* func) : mLibName(lib), mFuncName(func), mGPULib(nullptr), mGPUEntry(nullptr) {}
 

@@ -0,0 +1,105 @@
+# Copyright 2019-2026 CERN and copyright holders of ALICE O2.
+# See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
+# All rights not expressly granted are reserved.
+#
+# This software is distributed under the terms of the GNU General Public
+# License v3 (GPL Version 3), copied verbatim in the file "COPYING".
+#
+# In applying this license CERN does not waive the privileges and immunities
+# granted to it by virtue of its status as an Intergovernmental Organization
+# or submit itself to any jurisdiction.
+
+set(MODULE GPUTrackingCPUDevice)
+
+set(SRCS GPUReconstructionCPUDevice.cxx GPUReconstructionCPUDeviceGenRTC.cxx)
+set(HDRS GPUReconstructionCPUDevice.h GPUReconstructionCPUDeviceInternals.h GPUReconstructionCPUDeviceIncludesSystem.h)
+
+if(CMAKE_SYSTEM_NAME STREQUAL "Linux")
+  set(GPUCA_CPUDEVICE_RTC ON)
+else()
+  set(GPUCA_CPUDEVICE_RTC OFF)
+endif()
+message(STATUS "Building GPUTracking with CPUDevice support (RTC: ${GPUCA_CPUDEVICE_RTC})")
+
+if(ALIGPU_BUILD_TYPE STREQUAL "O2")
+  set(defineIncludeSrc "O2::${MODULE}")
+else()
+  set(defineIncludeSrc "${MODULE}")
+endif()
+
+if(GPUCA_CPUDEVICE_RTC) # RTC compiles the kernels with the host compiler and the host compiler flags, and loads them via dlopen
+  enable_language(ASM)
+  if(ALIGPU_BUILD_TYPE STREQUAL "Standalone")
+    get_filename_component(GPUDIR ${CMAKE_SOURCE_DIR}/../ ABSOLUTE)
+  else()
+    set(GPUDIR ${CMAKE_SOURCE_DIR}/GPU/GPUTracking)
+  endif()
+  string(TOUPPER "${CMAKE_BUILD_TYPE}" TMP_BUILD_TYPE_UPPER)
+  set(GPU_RTC_DEFINES "-D$<JOIN:$<TARGET_PROPERTY:${defineIncludeSrc},COMPILE_DEFINITIONS>,$<SEMICOLON>-D>")
+  set(GPU_RTC_INCLUDES "-I$<JOIN:$<FILTER:$<TARGET_PROPERTY:${defineIncludeSrc},INCLUDE_DIRECTORIES>,EXCLUDE,^/usr/include/?>,$<SEMICOLON>-I>"
+                       -I${CMAKE_SOURCE_DIR}/Detectors/Base/src
+                       -I${CMAKE_SOURCE_DIR}/Detectors/TRD/base/src
+  )
+  separate_arguments(GPU_RTC_FLAGS NATIVE_COMMAND "-std=c++${CMAKE_CXX_STANDARD} ${CMAKE_CXX_FLAGS} ${CMAKE_CXX_FLAGS_${TMP_BUILD_TYPE_UPPER}} -w")
+  set(GPU_RTC_LINK_FLAGS -fPIC -shared -fvisibility=hidden -fvisibility-inlines-hidden -static-libstdc++ -static-libgcc -Wl,--exclude-libs,ALL -Wl,-Bsymbolic -Wl,-z,defs) # Self-contained library, only the kernels are exported (version script at runtime)
+  set(GPU_RTC_SRC ${GPUDIR}/Base/cpudevice/GPUReconstructionCPUDevicertc.cxx)
+  set(GPU_RTC_BIN ${CMAKE_CURRENT_BINARY_DIR}/GPUReconstructionCPUDevicertc)
+
+  add_custom_command(
+      OUTPUT ${GPU_RTC_BIN}.src
+      COMMAND cp ${GPUDIR}/Base/GPUStdSystemHeaders.h ${GPU_RTC_BIN}.src
+      COMMAND cat ${GPUDIR}/Base/cpudevice/GPUReconstructionCPUDeviceIncludesSystem.h | grep -v GPUStdSystemHeaders.h >> ${GPU_RTC_BIN}.src
+      COMMAND ${CMAKE_CXX_COMPILER} ${GPU_RTC_DEFINES} ${GPU_RTC_INCLUDES} ${GPU_RTC_FLAGS} -x c++ -nostdinc -P -E -MD -MT ${GPU_RTC_BIN}.src -MF ${GPU_RTC_BIN}.src.d ${GPU_RTC_SRC} >> ${GPU_RTC_BIN}.src
+      DEPENDS ${GPU_RTC_SRC} ${GPUDIR}/Base/GPUStdSystemHeaders.h ${GPUDIR}/Base/cpudevice/GPUReconstructionCPUDeviceIncludesSystem.h
+      DEPFILE ${GPU_RTC_BIN}.src.d
+      COMMAND_EXPAND_LISTS
+      COMMENT "Preparing CPUDevice RTC source file ${GPU_RTC_BIN}.src"
+  )
+  add_custom_command(
+      OUTPUT ${GPU_RTC_BIN}.command
+      COMMAND echo -n "${CMAKE_CXX_COMPILER} ${GPU_RTC_FLAGS} ${GPU_RTC_LINK_FLAGS}" > ${GPU_RTC_BIN}.command
+      COMMAND_EXPAND_LISTS VERBATIM
+      COMMENT "Preparing CPUDevice RTC command file ${GPU_RTC_BIN}.command"
+  )
+  add_custom_command(
+      OUTPUT ${GPU_RTC_BIN}.command.no_fast_math
+      COMMAND echo -n "${GPUCA_CXX_NO_FAST_MATH_FLAGS}" > ${GPU_RTC_BIN}.command.no_fast_math
+      COMMAND_EXPAND_LISTS VERBATIM
+      COMMENT "Preparing CPUDevice RTC NO_FAST_MATH command file ${GPU_RTC_BIN}.command.no_fast_math"
+  )
+  foreach(TMP_RESOURCE src command command.no_fast_math)
+    create_binary_resource(${GPU_RTC_BIN}.${TMP_RESOURCE} ${GPU_RTC_BIN}.${TMP_RESOURCE}.o)
+    list(APPEND SRCS ${GPU_RTC_BIN}.${TMP_RESOURCE}.o)
+  endforeach()
+endif()
+
+if(ALIGPU_BUILD_TYPE STREQUAL "O2")
+  o2_add_library(
+    ${MODULE}
+    SOURCES ${SRCS}
+    PUBLIC_LINK_LIBRARIES O2::GPUTracking
+    PRIVATE_INCLUDE_DIRECTORIES
+      ${CMAKE_SOURCE_DIR}/Detectors/Base/src
+      ${CMAKE_SOURCE_DIR}/Detectors/TRD/base/src
+      ${CMAKE_SOURCE_DIR}/DataFormats/Reconstruction/src
+      ${CMAKE_CURRENT_SOURCE_DIR}
+    TARGETVARNAME targetName)
+endif()
+
+if(ALIGPU_BUILD_TYPE STREQUAL "Standalone")
+  set(targetName "${MODULE}")
+  add_library(${MODULE} SHARED ${SRCS})
+  add_library(O2::${MODULE} ALIAS ${MODULE})
+  target_link_libraries(${MODULE} PUBLIC O2::GPUTracking)
+  install(TARGETS ${MODULE})
+  include_directories(${CMAKE_CURRENT_SOURCE_DIR})
+endif()
+
+install(FILES ${HDRS} DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/GPU)
+
+target_compile_definitions(${targetName} PRIVATE $<TARGET_PROPERTY:O2::GPUTracking,COMPILE_DEFINITIONS>)
+target_link_libraries(${targetName} PRIVATE TBB::tbb)
+if(GPUCA_CPUDEVICE_RTC)
+  target_compile_definitions(${targetName} PRIVATE GPUCA_CPUDEVICE_RTC)
+  target_link_libraries(${targetName} PRIVATE ${CMAKE_DL_LIBS})
+endif()