From c44cd2540b8c4d5abfb007bb3bb83e15217076c7 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Sun, 11 Oct 2026 10:41:02 +0200 Subject: [PATCH] GPU: Add CPUDevice backend to emulate device-like processing on host --- GPU/Common/GPUCommonAlgorithm.h | 6 +- GPU/Common/GPUCommonArray.h | 2 +- GPU/Common/GPUCommonDef.h | 6 +- GPU/Common/GPUCommonDefAPI.h | 2 +- GPU/Common/GPUCommonMath.h | 2 +- GPU/GPUTracking/Base/GPUReconstruction.h | 2 +- ...ReconstructionAvailableBackends.template.h | 1 + .../Base/GPUReconstructionLibrary.cxx | 5 + GPU/GPUTracking/Base/cpudevice/CMakeLists.txt | 105 ++++ .../cpudevice/GPUReconstructionCPUDevice.cxx | 447 ++++++++++++++++++ .../cpudevice/GPUReconstructionCPUDevice.h | 81 ++++ .../GPUReconstructionCPUDeviceGenRTC.cxx | 367 ++++++++++++++ ...GPUReconstructionCPUDeviceIncludesSystem.h | 27 ++ .../GPUReconstructionCPUDeviceInternals.h | 88 ++++ .../GPUReconstructionCPUDevicertc.cxx | 28 ++ GPU/GPUTracking/CMakeLists.txt | 7 + .../DataTypes/GPUDataTypesConfig.h | 4 +- GPU/GPUTracking/Definitions/GPUSettingsList.h | 4 +- .../Definitions/clusterFinderDefs.h | 2 +- .../Global/GPUChainTrackingClusterizer.cxx | 11 +- GPU/GPUTracking/Standalone/cmake/config.cmake | 7 +- GPU/GPUTracking/TPCClusterFinder/CfUtils.h | 4 +- .../GPUTPCCFCheckPadBaseline.cxx | 2 +- .../TPCClusterFinder/GPUTPCCFClusterizer.inc | 2 +- .../GPUTPCCFNoiseSuppression.cxx | 2 +- 25 files changed, 1189 insertions(+), 25 deletions(-) create mode 100644 GPU/GPUTracking/Base/cpudevice/CMakeLists.txt create mode 100644 GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDevice.cxx create mode 100644 GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDevice.h create mode 100644 GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDeviceGenRTC.cxx create mode 100644 GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDeviceIncludesSystem.h create mode 100644 GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDeviceInternals.h create mode 100644 GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDevicertc.cxx diff --git a/GPU/Common/GPUCommonAlgorithm.h b/GPU/Common/GPUCommonAlgorithm.h index ef9b54ef3f96d..a153ae3c91f11 100644 --- a/GPU/Common/GPUCommonAlgorithm.h +++ b/GPU/Common/GPUCommonAlgorithm.h @@ -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 +#endif #define GPUCA_ALGORITHM_STD #endif @@ -280,7 +282,7 @@ GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end GPUCA_THREAD_INFO_ template 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) { diff --git a/GPU/Common/GPUCommonArray.h b/GPU/Common/GPUCommonArray.h index e647d8c953701..4a80cbe70cc4c 100644 --- a/GPU/Common/GPUCommonArray.h +++ b/GPU/Common/GPUCommonArray.h @@ -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 +#if !defined(GPUCA_GPUCODE_DEVICE) || defined(__CUDACC__) || defined(__HIPCC__) || defined(__CPUDEVICE__) // TODO: Get rid of GPUCommonArray once OpenCL supports #ifndef GPUCA_GPUCODE_COMPILEKERNELS #include #endif diff --git a/GPU/Common/GPUCommonDef.h b/GPU/Common/GPUCommonDef.h index 3b15e16314eef..5c68366702f93 100644 --- a/GPU/Common/GPUCommonDef.h +++ b/GPU/Common/GPUCommonDef.h @@ -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 diff --git a/GPU/Common/GPUCommonDefAPI.h b/GPU/Common/GPUCommonDefAPI.h index 9aa88f511f889..a294ea6e81b19 100644 --- a/GPU/Common/GPUCommonDefAPI.h +++ b/GPU/Common/GPUCommonDefAPI.h @@ -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 diff --git a/GPU/Common/GPUCommonMath.h b/GPU/Common/GPUCommonMath.h index 057d3afca0945..081af15347b10 100644 --- a/GPU/Common/GPUCommonMath.h +++ b/GPU/Common/GPUCommonMath.h @@ -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 diff --git a/GPU/GPUTracking/Base/GPUReconstruction.h b/GPU/GPUTracking/Base/GPUReconstruction.h index 993db1f381f17..eb2653fdfb66d 100644 --- a/GPU/GPUTracking/Base/GPUReconstruction.h +++ b/GPU/GPUTracking/Base/GPUReconstruction.h @@ -427,7 +427,7 @@ class GPUReconstruction void* mGPULib; void* mGPUEntry; }; - static std::shared_ptr sLibCUDA, sLibHIP, sLibOCL, sLibMETAL; + static std::shared_ptr sLibCUDA, sLibHIP, sLibOCL, sLibMETAL, sLibCPUDEVICE; // Debugging struct debugInternal; diff --git a/GPU/GPUTracking/Base/GPUReconstructionAvailableBackends.template.h b/GPU/GPUTracking/Base/GPUReconstructionAvailableBackends.template.h index 661a39e99b20f..93f6a1ae94bef 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionAvailableBackends.template.h +++ b/GPU/GPUTracking/Base/GPUReconstructionAvailableBackends.template.h @@ -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 diff --git a/GPU/GPUTracking/Base/GPUReconstructionLibrary.cxx b/GPU/GPUTracking/Base/GPUReconstructionLibrary.cxx index 104c91a6e94e3..614301cdeecba 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionLibrary.cxx +++ b/GPU/GPUTracking/Base/GPUReconstructionLibrary.cxx @@ -106,6 +106,10 @@ std::shared_ptr* 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::sLibCUDA(ne std::shared_ptr GPUReconstruction::sLibHIP(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingHIP" LIBRARY_EXTENSION, "GPUReconstruction_Create_HIP")); std::shared_ptr GPUReconstruction::sLibOCL(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingOCL" LIBRARY_EXTENSION, "GPUReconstruction_Create_OCL")); std::shared_ptr GPUReconstruction::sLibMETAL(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingMETAL" LIBRARY_EXTENSION, "GPUReconstruction_Create_METAL")); +std::shared_ptr 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) {} diff --git a/GPU/GPUTracking/Base/cpudevice/CMakeLists.txt b/GPU/GPUTracking/Base/cpudevice/CMakeLists.txt new file mode 100644 index 0000000000000..81926dda2753b --- /dev/null +++ b/GPU/GPUTracking/Base/cpudevice/CMakeLists.txt @@ -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$,$-D>") + set(GPU_RTC_INCLUDES "-I$,EXCLUDE,^/usr/include/?>,$-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_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() diff --git a/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDevice.cxx b/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDevice.cxx new file mode 100644 index 0000000000000..846dd32674be9 --- /dev/null +++ b/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDevice.cxx @@ -0,0 +1,447 @@ +// 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. + +/// \file GPUReconstructionCPUDevice.cxx +/// \author David Rohr + +#include "GPUReconstructionCPUDevice.h" +#include "GPUReconstructionCPUDeviceInternals.h" +#include "GPUReconstructionIncludes.h" +#include "GPUReconstructionKernelIncludes.h" +#include "GPUDef.h" +#include "GPUConstantMem.h" +#include "GPUDefParametersRuntime.h" +#include "GPULogging.h" + +#include +#include +#include +#include +#include + +#ifndef _WIN32 +#include +#endif + +using namespace o2::gpu; + +GPUReconstruction* GPUReconstruction_Create_CPUDEVICE(const GPUSettingsDeviceBackend& cfg) { return new GPUReconstructionCPUDevice(cfg); } + +GPUReconstructionCPUDevice::GPUReconstructionCPUDevice(const GPUSettingsDeviceBackend& cfg) : GPUReconstructionProcessing::KernelInterface(cfg, sizeof(GPUReconstructionDeviceBase)) +{ + if (mMaster == nullptr) { + mInternals = new GPUReconstructionCPUDeviceInternals; + *mParDevice = *mParCPU; // We run the CPU kernels on the "device", so we use the CPU kernel parameters + mParDevice->par_ALTERNATE_BORDER_SORT = 1; // Required by the GPU code path without RTC + mParDevice->par_CF_SCAN_WORKGROUP_SIZE = 1; // No block-wide scan with 1 thread per block, the clusterizer uses the CPU stream compaction + for (auto& v : mParDevice->par_LB_maxThreads) { + v = 1; + } + } + mDeviceBackendSettings->deviceType = DeviceType::CPUDEVICE; +} + +GPUReconstructionCPUDevice::~GPUReconstructionCPUDevice() +{ + Exit(); + if (mMaster == nullptr) { + delete mInternals; + } +} + +int32_t GPUReconstructionCPUDevice::GPUChkErrInternal(const int64_t retval, const char* file, int32_t line) const +{ + if (retval) { + GPUError("CPUDevice Error: %ld (%s:%d)", (int64_t)retval, file, line); + } + return retval != 0; +} + +int32_t GPUReconstructionCPUDevice::InitDevice_Runtime() +{ + if (GetProcessingSettings().rtc.enable && GetProcessingSettings().rtctech.runTest == 2) { + genAndLoadRTC(); + exit(0); + } + + if (mMaster == nullptr) { + mMultiprocessorCount = std::max(1, mMaxHostThreads); // One device block runs as one TBB task + mThreadCount = 1; + mWarpSize = 1; + mMaxBackendThreads = std::max(mMaxBackendThreads, mMultiprocessorCount); + mDeviceName = "CPUDevice"; + if (GetProcessingSettings().deviceTimers) { + if (GetProcessingSettings().debugLevel >= 2) { + GPUInfo("Disabling device timers for CPUDevice"); + } + mProcessingSettings->deviceTimers = 0; + } + + // The arena does not own threads, it limits the concurrency of the command queues and kernels to the configured number of host threads + mInternals->arena = std::make_unique(std::max(1, mMaxHostThreads)); + + mInternals->deviceMemory = alignedDefaultBufferAllocator(mDeviceMemorySize); + mInternals->hostMemory = alignedDefaultBufferAllocator(mHostMemorySize); + mInternals->constantMemory = alignedDefaultBufferAllocator(sizeof(GPUConstantMem)); + if ((mDeviceMemorySize && mInternals->deviceMemory == nullptr) || (mHostMemorySize && mInternals->hostMemory == nullptr) || mInternals->constantMemory == nullptr) { + GPUError("CPUDevice memory allocation error (trying %ld bytes device / %ld bytes host memory)", (int64_t)mDeviceMemorySize, (int64_t)mHostMemorySize); + return 1; + } + mDeviceMemoryBase = mInternals->deviceMemory; + mHostMemoryBase = mInternals->hostMemory; + mDeviceConstantMem = (GPUConstantMem*)mInternals->constantMemory; + if (GetProcessingSettings().debugLevel >= 1) { + GPUInfo("Memory ptrs: Device (%ld bytes): %p - Host (%ld bytes): %p", (int64_t)mDeviceMemorySize, mDeviceMemoryBase, (int64_t)mHostMemorySize, mHostMemoryBase); + memset(mHostMemoryBase, 0xDD, mHostMemorySize); + memset(mDeviceMemoryBase, 0xDD, mDeviceMemorySize); + } + + if (GetProcessingSettings().rtc.enable) { + genAndLoadRTC(); + } + + GPUInfo("CPUDevice Initialisation successfull (%d blocks, %d streams, %ld / %ld bytes host / device memory, Constant memory %ld)", mMultiprocessorCount, mNStreams, (int64_t)mHostMemorySize, (int64_t)mDeviceMemorySize, (int64_t)sizeof(GPUConstantMem)); + } else { + GPUReconstructionCPUDevice* master = dynamic_cast(mMaster); + mMultiprocessorCount = master->mMultiprocessorCount; + mThreadCount = master->mThreadCount; + mWarpSize = master->mWarpSize; + mMaxBackendThreads = master->mMaxBackendThreads; + mDeviceName = master->mDeviceName; + mDeviceConstantMem = master->mDeviceConstantMem; + mInternals = master->mInternals; + GPUInfo("CPUDevice Initialisation successfull (from master)"); + } + + mEventStorage.reserve(mEvents.size()); + for (uint32_t i = 0; i < mEvents.size(); i++) { + mEventStorage.emplace_back(mEvents[i].size()); + for (uint32_t j = 0; j < mEvents[i].size(); j++) { + mEvents[i][j].set(&mEventStorage[i][j]); + } + } + + return 0; +} + +int32_t GPUReconstructionCPUDevice::ExitDevice_Runtime() +{ + SynchronizeGPU(); + mEventStorage.clear(); + + if (mMaster == nullptr) { +#ifndef _WIN32 + for (auto* lib : mInternals->rtcLibraries) { + dlclose(lib); + } +#endif + mInternals->rtcLibraries.clear(); + mInternals->rtcKernels.clear(); + mInternals->arena.reset(); + alignedDefaultBufferDeleter()(mInternals->deviceMemory); + alignedDefaultBufferDeleter()(mInternals->hostMemory); + alignedDefaultBufferDeleter()(mInternals->constantMemory); + mInternals->deviceMemory = mInternals->hostMemory = mInternals->constantMemory = nullptr; + GPUInfo("CPUDevice Uninitialized"); + } + mDeviceMemoryBase = nullptr; + mHostMemoryBase = nullptr; + mDeviceConstantMem = nullptr; + + return 0; +} + +GPUReconstructionCPUDeviceEvent* GPUReconstructionCPUDevice::getEvent(deviceEvent& ev) +{ + return ev.get(); +} + +void GPUReconstructionCPUDeviceInternals::enqueue(int32_t stream, GPUReconstructionCPUDeviceTask&& task) +{ + auto& q = queues[stream]; + bool start = false; + { + std::lock_guard lock(q.mutex); + q.tasks.emplace_back(std::move(task)); + if (!q.active) { + q.active = start = true; + } + } + if (start) { + startQueue(stream); + } +} + +void GPUReconstructionCPUDeviceInternals::startQueue(int32_t stream) +{ + arena->enqueue([this, stream]() { processQueue(stream); }); +} + +void GPUReconstructionCPUDeviceInternals::processQueue(int32_t stream) +{ + auto& q = queues[stream]; + while (true) { + GPUReconstructionCPUDeviceTask* task; + { + std::lock_guard lock(q.mutex); + if (q.tasks.empty()) { + q.active = false; + q.cv.notify_all(); + return; + } + task = &q.tasks.front(); // Only this TBB task removes elements, and push_back does not invalidate references into a deque + } + while (task->nextWait < task->waits.size()) { + auto [e, value] = task->waits[task->nextWait]; + std::lock_guard lock(e->mutex); + if (e->nCompleted < value) { // Do not block the TBB thread, suspend the queue, completeEvent will start it again + e->waiting.push_back({value, this, stream}); + return; + } + task->nextWait++; + } + if (!error) { + try { + task->work(); + } catch (const std::exception& ex) { + GPUError("CPUDevice: Error in command queue %d: %s", stream, ex.what()); + error = 1; + } + } + if (task->record) { + completeEvent(task->record, task->recordValue); // Also after errors, so that nothing waits forever + } + std::lock_guard lock(q.mutex); + q.tasks.pop_front(); + } +} + +void GPUReconstructionCPUDeviceInternals::completeEvent(GPUReconstructionCPUDeviceEvent* e, uint64_t value) +{ + std::vector resume; + { + std::lock_guard lock(e->mutex); + e->nCompleted = std::max(e->nCompleted, value); + for (auto it = e->waiting.begin(); it != e->waiting.end();) { + if (it->value <= e->nCompleted) { + resume.emplace_back(*it); + it = e->waiting.erase(it); + } else { + it++; + } + } + e->cv.notify_all(); + } + for (auto& w : resume) { + w.internals->startQueue(w.stream); + } +} + +void GPUReconstructionCPUDevice::EnqueueTask(int32_t stream, deviceEvent* evList, int32_t nEvents, deviceEvent* ev, std::function&& work) +{ + if (stream < 0 || stream >= mNStreams) { + throw std::runtime_error("CPUDevice: invalid stream " + std::to_string(stream)); + } + GPUReconstructionCPUDeviceTask task; + if (evList) { // Capture the event states at enqueue time, such that the queue waits for the markers recorded so far, like cudaStreamWaitEvent + for (int32_t k = 0; k < nEvents; k++) { + auto* e = getEvent(evList[k]); + std::lock_guard lock(e->mutex); + task.waits.emplace_back(e, e->nRecorded); + } + } + if (ev) { + task.record = getEvent(*ev); + std::lock_guard lock(task.record->mutex); + task.recordValue = ++task.record->nRecorded; + } + task.work = std::move(work); + mInternals->enqueue(stream, std::move(task)); +} + +void GPUReconstructionCPUDevice::SynchronizeStream(int32_t stream) +{ + auto& q = mInternals->queues[stream]; + std::unique_lock lock(q.mutex); + q.cv.wait(lock, [&q]() { return !q.active; }); +} + +void GPUReconstructionCPUDevice::SynchronizeGPU() +{ + for (uint32_t i = 0; i < mInternals->queues.size(); i++) { + SynchronizeStream(i); + } +} + +void GPUReconstructionCPUDevice::SynchronizeEvents(deviceEvent* evList, int32_t nEvents) +{ + for (int32_t i = 0; i < nEvents; i++) { + auto* e = getEvent(evList[i]); + std::unique_lock lock(e->mutex); + const uint64_t value = e->nRecorded; + e->cv.wait(lock, [e, value]() { return e->nCompleted >= value; }); + } +} + +void GPUReconstructionCPUDevice::StreamWaitForEvents(int32_t stream, deviceEvent* evList, int32_t nEvents) +{ + if (nEvents) { + EnqueueTask(stream, evList, nEvents, nullptr, []() {}); + } +} + +bool GPUReconstructionCPUDevice::IsEventDone(deviceEvent* evList, int32_t nEvents) +{ + for (int32_t i = 0; i < nEvents; i++) { + auto* e = getEvent(evList[i]); + std::lock_guard lock(e->mutex); + if (e->nCompleted < e->nRecorded) { + return false; + } + } + return true; +} + +void GPUReconstructionCPUDevice::RecordMarker(deviceEvent* ev, int32_t stream) +{ + EnqueueTask(stream, nullptr, 0, ev, []() {}); +} + +int32_t GPUReconstructionCPUDevice::GPUDebug(const char* state, int32_t stream, bool force) +{ + if (mInternals->error) { + GPUError("CPUDevice Error while running (%s) (Stream %d)", state, stream); + return 1; + } + if (!force && GetProcessingSettings().debugLevel <= 0) { + return 0; + } + if (stream == -1) { + SynchronizeGPU(); + } else { + SynchronizeStream(stream); + } + if (mInternals->error) { + GPUError("CPUDevice Error while synchronizing (%s) (Stream %d)", state, stream); + return 1; + } + if (GetProcessingSettings().debugLevel >= 3) { + GPUInfo("GPU Sync Done"); + } + return 0; +} + +size_t GPUReconstructionCPUDevice::GPUMemCpy(void* dst, const void* src, size_t size, int32_t stream, int32_t toGPU, deviceEvent* ev, deviceEvent* evList, int32_t nEvents) +{ + if (GetProcessingSettings().debugLevel >= 3) { + stream = -1; + } + if (stream == -1) { + SynchronizeGPU(); + if (size) { + memcpy(dst, src, size); + } + if (ev) { + auto* e = getEvent(*ev); + uint64_t value; + { + std::lock_guard lock(e->mutex); + value = ++e->nRecorded; + } + mInternals->completeEvent(e, value); + } + } else { + EnqueueTask(stream, evList, evList ? nEvents : 0, ev, [dst, src, size]() { + if (size) { + memcpy(dst, src, size); + } + }); + } + if (GetProcessingSettings().serializeGPU & 2) { + GPUDebug(("GPUMemCpy " + std::to_string(toGPU)).c_str(), stream, true); + } + return size; +} + +size_t GPUReconstructionCPUDevice::WriteToConstantMemory(size_t offset, const void* src, size_t size, int32_t stream, deviceEvent* ev) +{ + char* dst = (char*)mDeviceConstantMem + offset; + if (stream == -1) { + SynchronizeGPU(); + memcpy(dst, src, size); + } else { + // The source may be a temporary, the CUDA backend stages the copy too, so we take a copy for the asynchronous transfer + auto buffer = std::make_shared>((const char*)src, (const char*)src + size); + EnqueueTask(stream, nullptr, 0, ev, [dst, buffer]() { memcpy(dst, buffer->data(), buffer->size()); }); + } + if (GetProcessingSettings().serializeGPU & 2) { + GPUDebug("WriteToConstantMemory", stream, true); + } + return size; +} + +template +inline void GPUReconstructionCPUDevice::runKernelBackend(const krnlSetupTime& _xyz, const Args&... args) +{ + auto& x = _xyz.x; + auto& y = _xyz.y; + auto& z = _xyz.z; + if (x.nThreads != 1) { + throw std::runtime_error("CPUDevice backend supports only kernels with nThreads = 1"); + } + const uint32_t nBlocks = x.nBlocks; + const int32_t iSector = y.index; + GPUConstantMem* constantMem = mDeviceConstantMem; + using rtcKernel = void (*)(GPUConstantMem*, int32_t, uint32_t, uint32_t, uint32_t, Args...); + rtcKernel rtc = mInternals->rtcKernels.size() ? reinterpret_cast(mInternals->rtcKernels[GetKernelNum()]) : nullptr; + + // The kernel arguments are captured by value, since the kernel runs asynchronously in the command queue + EnqueueTask(x.stream, z.evList, z.evList ? z.nEvents : 0, z.ev, [nBlocks, iSector, constantMem, rtc, ... args = args]() { + tbb::this_task_arena::isolate([&] { // Do not pick up tasks of other command queues while waiting inside parallel_for + tbb::parallel_for(tbb::blocked_range(0, nBlocks, 1), [&](const tbb::blocked_range& r) { + if (rtc) { + rtc(constantMem, iSector, nBlocks, r.begin(), r.end(), args...); + } else { + for (uint32_t iBlock = r.begin(); iBlock < r.end(); iBlock++) { + typename T::GPUSharedMemory smem; + T::template Thread(nBlocks, 1, iBlock, 0, smem, T::Processor(*constantMem)[iSector], args...); + } + } + }); + }); + }); +} + +template <> +inline void GPUReconstructionCPUDevice::runKernelBackend(const krnlSetupTime& _xyz, void* const& ptr, uint64_t const& size) +{ + EnqueueTask(_xyz.x.stream, _xyz.z.evList, _xyz.z.evList ? _xyz.z.nEvents : 0, _xyz.z.ev, [ptr, size]() { memset(ptr, 0, size); }); +} + +#include "GPUReconstructionProcessingKernels.inc" +template void GPUReconstructionProcessing::KernelInterface::runKernelVirtual(const int num, const void* args); + +// Entry point of an RTC-compiled kernel, running the blocks [iBlockStart, iBlockEnd) with one thread per block. Must match the rtcKernel function pointer type in GPUReconstructionCPUDevice.cxx. +#define GPUCA_CPUDEVICE_KRNL_RTC(x_class, x_attributes, x_arguments, x_forward, x_types, ...) \ + extern "C" __attribute__((visibility("default"))) void GPUCA_M_CAT(krnl_, GPUCA_M_KRNL_NAME(x_class))(GPUConstantMem * pConstant, int32_t _iSector_internal, uint32_t nBlocks, uint32_t iBlockStart, uint32_t iBlockEnd GPUCA_M_STRIP(x_arguments)) \ + { \ + for (uint32_t iBlock = iBlockStart; iBlock < iBlockEnd; iBlock++) { \ + typename GPUCA_M_STRIP_FIRST(x_class)::GPUSharedMemory smem; \ + GPUCA_M_STRIP_FIRST(x_class)::template Thread(nBlocks, 1, iBlock, 0, smem, GPUCA_M_STRIP_FIRST(x_class)::Processor(*pConstant)[_iSector_internal] GPUCA_M_STRIP(x_forward)); \ + } \ + } + +void GPUReconstructionCPUDevice::getRTCKernelCalls(std::vector& kernels) +{ +#define GPUCA_KRNL(...) kernels.emplace_back(GPUCA_M_STR(GPUCA_CPUDEVICE_KRNL_RTC(__VA_ARGS__))); +#include "GPUReconstructionKernelList.h" +#undef GPUCA_KRNL +} diff --git a/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDevice.h b/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDevice.h new file mode 100644 index 0000000000000..239afd60c6d0e --- /dev/null +++ b/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDevice.h @@ -0,0 +1,81 @@ +// 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. + +/// \file GPUReconstructionCPUDevice.h +/// \author David Rohr + +#ifndef GPURECONSTRUCTIONCPUDEVICE_H +#define GPURECONSTRUCTIONCPUDEVICE_H + +#include "GPUReconstructionDeviceBase.h" +#include +#include +#include + +#ifdef _WIN32 +extern "C" __declspec(dllexport) o2::gpu::GPUReconstruction* GPUReconstruction_Create_CPUDEVICE(const o2::gpu::GPUSettingsDeviceBackend& cfg); +#else +extern "C" o2::gpu::GPUReconstruction* GPUReconstruction_Create_CPUDEVICE(const o2::gpu::GPUSettingsDeviceBackend& cfg); +#endif + +namespace o2::gpu +{ +struct GPUReconstructionCPUDeviceInternals; +struct GPUReconstructionCPUDeviceEvent; + +// Backend which runs on the host, but behaves like a GPU device: It has a separate device memory region with explicit memcopies, +// and multiple in-order command queues (streams) which are processed asynchronously by TBB. One device block is one TBB task, +// the number of device threads per block is fixed to 1, such that the normal CPU kernels can be used. +class GPUReconstructionCPUDevice : public GPUReconstructionProcessing::KernelInterface +{ + public: + GPUReconstructionCPUDevice(const GPUSettingsDeviceBackend& cfg); + ~GPUReconstructionCPUDevice() override; + + template + void runKernelBackend(const krnlSetupTime& _xyz, const Args&... args); + + protected: + int32_t InitDevice_Runtime() override; + int32_t ExitDevice_Runtime() override; + + int32_t GPUChkErrInternal(const int64_t retval, const char* file, int32_t line) const override; + + void SynchronizeGPU() override; + int32_t GPUDebug(const char* state = "UNKNOWN", int32_t stream = -1, bool force = false) override; + void SynchronizeStream(int32_t stream) override; + void SynchronizeEvents(deviceEvent* evList, int32_t nEvents = 1) override; + void StreamWaitForEvents(int32_t stream, deviceEvent* evList, int32_t nEvents = 1) override; + bool IsEventDone(deviceEvent* evList, int32_t nEvents = 1) override; + int32_t registerMemoryForGPU_internal(const void* ptr, size_t size) override { return 0; } + int32_t unregisterMemoryForGPU_internal(const void* ptr) override { return 0; } + + size_t WriteToConstantMemory(size_t offset, const void* src, size_t size, int32_t stream = -1, deviceEvent* ev = nullptr) override; + size_t GPUMemCpy(void* dst, const void* src, size_t size, int32_t stream, int32_t toGPU, deviceEvent* ev = nullptr, deviceEvent* evList = nullptr, int32_t nEvents = 1) override; + void ReleaseEvent(deviceEvent ev) override {} + void RecordMarker(deviceEvent* ev, int32_t stream) override; + + GPUReconstructionCPUDeviceInternals* mInternals = nullptr; + + private: + // Enqueue a task into the command queue of the stream, which waits for evList, runs work, and then marks ev as completed + void EnqueueTask(int32_t stream, deviceEvent* evList, int32_t nEvents, deviceEvent* ev, std::function&& work); + static GPUReconstructionCPUDeviceEvent* getEvent(deviceEvent& ev); + std::vector> mEventStorage; // Backs the deviceEvents in mEvents, inner vectors are never resized + + int32_t genRTC(std::string& filename, uint32_t& nCompile); + void getRTCKernelCalls(std::vector& kernels); + void genAndLoadRTC(); +}; + +} // namespace o2::gpu + +#endif diff --git a/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDeviceGenRTC.cxx b/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDeviceGenRTC.cxx new file mode 100644 index 0000000000000..14501ffc9da6b --- /dev/null +++ b/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDeviceGenRTC.cxx @@ -0,0 +1,367 @@ +// 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. + +/// \file GPUReconstructionCPUDeviceGenRTC.cxx +/// \author David Rohr + +#include "GPUReconstructionCPUDevice.h" +#include "GPULogging.h" + +#include +#include + +#ifdef GPUCA_CPUDEVICE_RTC + +#include "GPUReconstructionCPUDeviceInternals.h" +#include "GPUReconstructionKernelIncludes.h" +#include "GPUReconstructionKernelMacros.h" +#include "GPUParamRTC.h" +#include "GPUDefParametersLoad.inc" +#include "GPUNoFastMathKernels.h" +#include "Framework/SHA1.h" +#include "utils/qGetLdBinarySymbols.h" +#include "utils/timer.h" + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +QGET_LD_BINARY_SYMBOLS(GPUReconstructionCPUDevicertc_src); +QGET_LD_BINARY_SYMBOLS(GPUReconstructionCPUDevicertc_command); +QGET_LD_BINARY_SYMBOLS(GPUReconstructionCPUDevicertc_command_no_fast_math); + +using namespace o2::gpu; + +static constexpr const char* RTC_SRC_EXTENSION = ".cxx"; +static constexpr const char* RTC_BIN_EXTENSION = ".so"; +static constexpr const char* RTC_MAP_EXTENSION = ".map"; + +static bool kernelWithoutRTC(const std::string& name) { return name == "GPUTPCCFCheckPadBaseline"; } // Device code requires more than 1 thread per block, use the host-compiled kernel + +int32_t GPUReconstructionCPUDevice::genRTC(std::string& filename, uint32_t& nCompile) +{ + const std::string_view rtcSource(_binary_GPUReconstructionCPUDevicertc_src_start, _binary_GPUReconstructionCPUDevicertc_src_len); + const std::string rtcparam = std::string("#define GPUCA_RTC_CODE\n") + + std::string(GetProcessingSettings().rtc.optSpecialCode ? "#define GPUCA_RTC_SPECIAL_CODE(...) __VA_ARGS__\n" : "#define GPUCA_RTC_SPECIAL_CODE(...)\n") + + std::string(GetProcessingSettings().rtc.optConstexpr ? "#define GPUCA_RTC_CONSTEXPR constexpr\n" : "#define GPUCA_RTC_CONSTEXPR\n") + + GPUParamRTC::generateRTCCode(param(), GetProcessingSettings().rtc.optConstexpr); + if (filename == "") { + filename = "/tmp/o2cagpu_rtc_cpudevice_"; + } + filename += std::to_string(getpid()); + filename += "_"; + filename += std::to_string(rand()); + + std::vector kernels; + getRTCKernelCalls(kernels); + for (uint32_t i = 0; i < kernels.size(); i++) { + if (kernelWithoutRTC(GetKernelName(i))) { + kernels[i].clear(); + } + } + std::string kernelsall; + for (uint32_t i = 0; i < kernels.size(); i++) { + kernelsall += kernels[i] + "\n"; + } + + std::string baseCommand = (GetProcessingSettings().rtctech.prependCommand != "" ? (GetProcessingSettings().rtctech.prependCommand + " ") : ""); + baseCommand += (getenv("O2_GPU_RTC_OVERRIDE_CMD") ? std::string(getenv("O2_GPU_RTC_OVERRIDE_CMD")) : std::string(_binary_GPUReconstructionCPUDevicertc_command_start, _binary_GPUReconstructionCPUDevicertc_command_len)); + baseCommand += std::string(" ") + GetProcessingSettings().rtctech.overrideArchitecture; + + if (GetProcessingSettings().rtctech.loadLaunchBoundsFromFile.size()) { + FILE* fp = fopen(GetProcessingSettings().rtctech.loadLaunchBoundsFromFile.c_str(), "rb"); + if (fp == nullptr) { + throw std::runtime_error("Cannot open launch bounds parameter module file"); + } + fseek(fp, 0, SEEK_END); + size_t size = ftell(fp); + if (size != sizeof(*mParDevice)) { + throw std::runtime_error("launch bounds parameter file has incorrect size"); + } + fseek(fp, 0, SEEK_SET); + if (fread(mParDevice, 1, size, fp) != size) { + throw std::runtime_error("Error reading launch bounds parameter file"); + } + fclose(fp); + } + const std::string launchBounds = o2::gpu::internal::GPUDefParametersExport(*mParDevice, true, 0) + "#define GPUCA_WARP_SIZE " + std::to_string(mWarpSize) + "\n"; + if (GetProcessingSettings().rtctech.printLaunchBounds || GetProcessingSettings().debugLevel >= 3) { + GPUInfo("RTC Launch Bounds:\n%s", launchBounds.c_str()); + if (GetProcessingSettings().rtctech.printLaunchBounds >= 2) { + return 1; + } + } + + const std::string compilerVersions = getBackendVersions(); + + char shasource[21], shaparam[21], shacmd[21], shakernels[21], shabounds[21], shaversion[21]; + if (GetProcessingSettings().rtc.cacheOutput) { + o2::framework::internal::SHA1(shasource, rtcSource.data(), rtcSource.size()); + o2::framework::internal::SHA1(shaparam, rtcparam.c_str(), rtcparam.size()); + o2::framework::internal::SHA1(shacmd, baseCommand.c_str(), baseCommand.size()); + o2::framework::internal::SHA1(shakernels, kernelsall.c_str(), kernelsall.size()); + o2::framework::internal::SHA1(shabounds, launchBounds.c_str(), launchBounds.size()); + o2::framework::internal::SHA1(shaversion, compilerVersions.c_str(), compilerVersions.size()); + } + + nCompile = GetProcessingSettings().rtc.compilePerKernel ? kernels.size() : 1; + bool cacheLoaded = false; + int32_t fd = 0; + const std::string cacheFile = GetProcessingSettings().rtctech.cacheFolder + "/rtc.cpudevice.cache"; + if (GetProcessingSettings().rtc.cacheOutput) { + if (GetProcessingSettings().rtctech.cacheFolder != ".") { + std::filesystem::create_directories(GetProcessingSettings().rtctech.cacheFolder); + } + if (GetProcessingSettings().rtctech.cacheMutex) { + mode_t mask = S_IRUSR | S_IWUSR | S_IRGRP | S_IWGRP | S_IROTH | S_IWOTH; + fd = open((GetProcessingSettings().rtctech.cacheFolder + "/cache.lock").c_str(), O_RDWR | O_CREAT | O_CLOEXEC, mask); + if (fd == -1) { + throw std::runtime_error("Error opening rtc cache mutex lock file"); + } + fchmod(fd, mask); + if (lockf(fd, F_LOCK, 0)) { + throw std::runtime_error("Error locking rtc cache mutex file"); + } + } + + FILE* fp = fopen(cacheFile.c_str(), "rb"); + char sharead[20]; + if (fp) { + size_t len; + while (true) { + auto checkSHA = [&](const char* shacmp, const char* name) { + if (fread(sharead, 1, 20, fp) != 20) { + throw std::runtime_error("Cache file corrupt"); + } + if (!GetProcessingSettings().rtctech.ignoreCacheValid && memcmp(sharead, shacmp, 20)) { + GPUInfo("Cache file content outdated (%s)", name); + return 1; + } + return 0; + }; + if (checkSHA(shasource, "source") || + checkSHA(shaparam, "param") || + checkSHA(shacmd, "command line") || + checkSHA(shakernels, "kernel definitions") || + checkSHA(shabounds, "launch bounds") || + checkSHA(shaversion, "compiler versions")) { + break; + } + GPUSettingsProcessingRTC cachedSettings; + static_assert(std::is_trivially_copyable_v == true, "GPUSettingsProcessingRTC must be POD"); + if (fread(&cachedSettings, sizeof(cachedSettings), 1, fp) != 1) { + throw std::runtime_error("Cache file corrupt"); + } + if (!GetProcessingSettings().rtctech.ignoreCacheValid && !(cachedSettings == GetProcessingSettings().rtc)) { + GPUInfo("Cache file content outdated (rtc parameters)"); + break; + } + std::vector buffer; + for (uint32_t i = 0; i < nCompile; i++) { + if (fread(&len, sizeof(len), 1, fp) != 1) { + throw std::runtime_error("Cache file corrupt"); + } + buffer.resize(len); + if (fread(buffer.data(), 1, len, fp) != len) { + throw std::runtime_error("Cache file corrupt"); + } + FILE* fp2 = fopen((filename + "_" + std::to_string(i) + RTC_BIN_EXTENSION).c_str(), "w+b"); + if (fp2 == nullptr) { + throw std::runtime_error("Cannot open tmp file"); + } + if (fwrite(buffer.data(), 1, len, fp2) != len) { + throw std::runtime_error("Error writing file"); + } + fclose(fp2); + } + GPUInfo("Using RTC cache file"); + cacheLoaded = true; + break; + }; + fclose(fp); + } + } + if (!cacheLoaded) { + if (GetProcessingSettings().debugLevel >= 0) { + GPUInfo("Starting CPUDevice RTC Compilation"); + } + HighResTimer rtcTimer; + rtcTimer.ResetStart(); + tbb::parallel_for(0, nCompile, [&](auto i) { + const std::string fileBase = filename + "_" + std::to_string(i); + if (GetProcessingSettings().debugLevel >= 3) { + printf("Compiling %s\n", (fileBase + RTC_SRC_EXTENSION).c_str()); + } + FILE* fp = fopen((fileBase + RTC_SRC_EXTENSION).c_str(), "w+b"); + if (fp == nullptr) { + throw std::runtime_error("Error opening file"); + } + + const std::string& kernel = GetProcessingSettings().rtc.compilePerKernel ? kernels[i] : kernelsall; + + bool deterministic = GetProcessingSettings().rtc.deterministic || (GetProcessingSettings().rtc.compilePerKernel && o2::gpu::internal::noFastMathKernels.find(GetKernelName(i)) != o2::gpu::internal::noFastMathKernels.end()); + const std::string deterministicStr = std::string(deterministic ? "#define GPUCA_DETERMINISTIC_CODE(det, indet) det\n" : "#define GPUCA_DETERMINISTIC_CODE(det, indet) indet\n"); + + if (fwrite(deterministicStr.c_str(), 1, deterministicStr.size(), fp) != deterministicStr.size() || + fwrite(rtcparam.c_str(), 1, rtcparam.size(), fp) != rtcparam.size() || + fwrite(launchBounds.c_str(), 1, launchBounds.size(), fp) != launchBounds.size() || + fwrite(rtcSource.data(), 1, rtcSource.size(), fp) != rtcSource.size() || + fwrite(kernel.c_str(), 1, kernel.size(), fp) != kernel.size()) { + throw std::runtime_error("Error writing file"); + } + fclose(fp); + + // Export only the kernel entry points, everything else stays local to the RTC library to avoid symbol collisions + fp = fopen((fileBase + RTC_MAP_EXTENSION).c_str(), "w+b"); + if (fp == nullptr || fprintf(fp, "{ global: krnl_*; local: *; };\n") < 0) { + throw std::runtime_error("Error writing version script"); + } + fclose(fp); + + std::string command = baseCommand; + if (deterministic) { + command += std::string(" ") + std::string(_binary_GPUReconstructionCPUDevicertc_command_no_fast_math_start, _binary_GPUReconstructionCPUDevicertc_command_no_fast_math_len); + } + command += " -Wl,--version-script=" + fileBase + RTC_MAP_EXTENSION + " -x c++ " + fileBase + RTC_SRC_EXTENSION + " -o " + fileBase + RTC_BIN_EXTENSION; + if (GetProcessingSettings().debugLevel < 0) { + command += " &> /dev/null"; + } else if (GetProcessingSettings().debugLevel < 2) { + command += " > /dev/null"; + } + if (GetProcessingSettings().debugLevel >= 3) { + printf("Running command %s\n", command.c_str()); + } + if (system(command.c_str())) { + if (GetProcessingSettings().debugLevel >= 3) { + printf("Source code file: %s", (fileBase + RTC_SRC_EXTENSION).c_str()); + } + throw std::runtime_error("Error during CPUDevice RTC compilation"); + } // clang-format off + }, tbb::simple_partitioner()); // clang-format on + if (GetProcessingSettings().debugLevel >= 0) { + GPUInfo("RTC Compilation finished (%f seconds)", rtcTimer.GetCurrentElapsedTime()); + } + if (GetProcessingSettings().rtc.cacheOutput) { + FILE* fp = fopen(cacheFile.c_str(), "w+b"); + if (fp == nullptr) { + throw std::runtime_error("Cannot open cache file for writing"); + } + GPUInfo("Storing RTC compilation result in cache file"); + + if (fwrite(shasource, 1, 20, fp) != 20 || + fwrite(shaparam, 1, 20, fp) != 20 || + fwrite(shacmd, 1, 20, fp) != 20 || + fwrite(shakernels, 1, 20, fp) != 20 || + fwrite(shabounds, 1, 20, fp) != 20 || + fwrite(shaversion, 1, 20, fp) != 20 || + fwrite(&GetProcessingSettings().rtc, sizeof(GetProcessingSettings().rtc), 1, fp) != 1) { + throw std::runtime_error("Error writing cache file"); + } + + std::vector buffer; + for (uint32_t i = 0; i < nCompile; i++) { + FILE* fp2 = fopen((filename + "_" + std::to_string(i) + RTC_BIN_EXTENSION).c_str(), "rb"); + if (fp2 == nullptr) { + throw std::runtime_error("Cannot open RTC library file"); + } + fseek(fp2, 0, SEEK_END); + size_t size = ftell(fp2); + buffer.resize(size); + fseek(fp2, 0, SEEK_SET); + if (fread(buffer.data(), 1, size, fp2) != size) { + throw std::runtime_error("Error reading RTC library file"); + } + fclose(fp2); + + if (fwrite(&size, sizeof(size), 1, fp) != 1 || + fwrite(buffer.data(), 1, size, fp) != size) { + throw std::runtime_error("Error writing cache file"); + } + } + fclose(fp); + } + } + if (GetProcessingSettings().rtc.cacheOutput && GetProcessingSettings().rtctech.cacheMutex) { + if (lockf(fd, F_ULOCK, 0)) { + throw std::runtime_error("Error unlocking RTC cache mutex file"); + } + close(fd); + } + + return 0; +} + +void GPUReconstructionCPUDevice::genAndLoadRTC() +{ + std::string filename = ""; + uint32_t nCompile = 0; + if (genRTC(filename, nCompile)) { + throw std::runtime_error("Runtime compilation failed"); + } + for (uint32_t i = 0; i < nCompile; i++) { + const std::string fileBase = filename + "_" + std::to_string(i); + if (GetProcessingSettings().rtctech.runTest != 2) { + void* lib = dlopen((fileBase + RTC_BIN_EXTENSION).c_str(), RTLD_NOW | RTLD_LOCAL); + if (lib == nullptr) { + throw std::runtime_error(std::string("Error loading CPUDevice RTC library: ") + dlerror()); + } + mInternals->rtcLibraries.emplace_back(lib); + } + if (!GetProcessingSettings().rtctech.keepTempFiles) { + remove((fileBase + RTC_SRC_EXTENSION).c_str()); + remove((fileBase + RTC_BIN_EXTENSION).c_str()); + remove((fileBase + RTC_MAP_EXTENSION).c_str()); + } + } + if (GetProcessingSettings().rtctech.runTest == 2) { + return; + } + + const bool perKernel = GetProcessingSettings().rtc.compilePerKernel; + mInternals->rtcKernels.resize(GetNKernels()); + uint32_t j = 0; +#define GPUCA_KRNL(x_class, ...) \ + if (GetKernelNum() != j) { \ + GPUFatal("kernel numbers out of sync"); \ + } \ + mInternals->rtcKernels[j] = dlsym(mInternals->rtcLibraries[perKernel ? j : 0], GPUCA_M_STR(GPUCA_M_CAT(krnl_, GPUCA_M_KRNL_NAME(x_class)))); \ + if (mInternals->rtcKernels[j] == nullptr && !kernelWithoutRTC(GetKernelName(j))) { \ + GPUFatal("Error loading RTC kernel %s", GPUCA_M_STR(GPUCA_M_CAT(krnl_, GPUCA_M_KRNL_NAME(x_class)))); \ + } \ + j++; +#include "GPUReconstructionKernelList.h" +#undef GPUCA_KRNL + if (GetProcessingSettings().debugLevel >= 1) { + GPUInfo("Loaded %u CPUDevice RTC kernels from %u libraries", j, (uint32_t)mInternals->rtcLibraries.size()); + } +} + +#else // GPUCA_CPUDEVICE_RTC + +using namespace o2::gpu; + +int32_t GPUReconstructionCPUDevice::genRTC(std::string& filename, uint32_t& nCompile) +{ + throw std::runtime_error("CPUDevice backend compiled without RTC support (requires Linux)"); +} + +void GPUReconstructionCPUDevice::genAndLoadRTC() +{ + throw std::runtime_error("CPUDevice backend compiled without RTC support (requires Linux)"); +} + +#endif // GPUCA_CPUDEVICE_RTC diff --git a/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDeviceIncludesSystem.h b/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDeviceIncludesSystem.h new file mode 100644 index 0000000000000..68fa8a5372100 --- /dev/null +++ b/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDeviceIncludesSystem.h @@ -0,0 +1,27 @@ +// 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. + +/// \file GPUReconstructionCPUDeviceIncludesSystem.h +/// \author David Rohr + +#ifndef O2_GPU_GPURECONSTRUCTIONCPUDEVICEINCLUDESSYSTEM_H +#define O2_GPU_GPURECONSTRUCTIONCPUDEVICEINCLUDESSYSTEM_H + +#ifndef GPUCA_GPUCODE_GENRTC +#include "GPUStdSystemHeaders.h" +#endif + +#include +#include +#include +#include + +#endif // O2_GPU_GPURECONSTRUCTIONCPUDEVICEINCLUDESSYSTEM_H diff --git a/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDeviceInternals.h b/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDeviceInternals.h new file mode 100644 index 0000000000000..1499fa98847a3 --- /dev/null +++ b/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDeviceInternals.h @@ -0,0 +1,88 @@ +// 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. + +/// \file GPUReconstructionCPUDeviceInternals.h +/// \author David Rohr + +#ifndef GPURECONSTRUCTIONCPUDEVICEINTERNALS_H +#define GPURECONSTRUCTIONCPUDEVICEINTERNALS_H + +#include "GPUReconstructionCPUDevice.h" + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +namespace o2::gpu +{ +struct GPUReconstructionCPUDeviceInternals; + +// Event of the CPUDevice backend. Recording a marker increments nRecorded when the marker is enqueued, the command queue sets nCompleted when the marker is processed. A wait captures the value of nRecorded at the time the wait is enqueued, so it behaves like cudaStreamWaitEvent. +struct GPUReconstructionCPUDeviceEvent { + struct waitingQueue { + uint64_t value; + GPUReconstructionCPUDeviceInternals* internals; + int32_t stream; + }; + std::mutex mutex; + std::condition_variable cv; // For host threads waiting for the event + std::vector waiting; // Suspended command queues, which are resumed when the event reaches the value + uint64_t nRecorded = 0; + uint64_t nCompleted = 0; +}; + +// Command in a command queue: wait for events, run the work, mark an event as completed +struct GPUReconstructionCPUDeviceTask { + std::vector> waits; + uint32_t nextWait = 0; + std::function work; + GPUReconstructionCPUDeviceEvent* record = nullptr; + uint64_t recordValue = 0; +}; + +// In-order command queue. It does not own a thread: When commands are enqueued, a TBB task is started which processes them, and which returns once the queue is empty, or when the next command must wait for an event. In the latter case, the queue is registered at the event, +// and a new TBB task is started when the event completes. At most one TBB task processes a queue at a time. +struct GPUReconstructionCPUDeviceQueue { + std::mutex mutex; + std::condition_variable cv; // For host threads synchronizing the queue + std::deque tasks; + bool active = false; // Queue has pending commands, and is either being processed by a TBB task or suspended waiting for an event +}; + +struct GPUReconstructionCPUDeviceInternals { + std::unique_ptr arena; + std::array queues; + std::atomic error{0}; // Set if a command in a command queue has failed + + char* deviceMemory = nullptr; // "Device" memory, a separate host memory region + char* hostMemory = nullptr; // Host memory, corresponds to page-locked host memory of real GPUs + char* constantMemory = nullptr; // "Device" constant memory, holding the device copy of GPUConstantMem + + std::vector rtcLibraries; // dlopen handles of RTC-compiled kernel libraries + std::vector rtcKernels; // RTC kernel entry points, indexed by kernel number + + void enqueue(int32_t stream, GPUReconstructionCPUDeviceTask&& task); + void startQueue(int32_t stream); + void processQueue(int32_t stream); + void completeEvent(GPUReconstructionCPUDeviceEvent* e, uint64_t value); +}; + +} // namespace o2::gpu + +#endif diff --git a/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDevicertc.cxx b/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDevicertc.cxx new file mode 100644 index 0000000000000..5cabdc1b23160 --- /dev/null +++ b/GPU/GPUTracking/Base/cpudevice/GPUReconstructionCPUDevicertc.cxx @@ -0,0 +1,28 @@ +// 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. + +/// \file GPUReconstructionCPUDevicertc.cxx +/// \author David Rohr + +#define __CPUDEVICE__ // Compile as device code, which runs on the host +#define GPUCA_DEVICE_SINGLE_THREAD // Device code with 1 thread per block +#define GPUCA_GPUCODE_GENRTC +#define GPUCA_GPUCODE_COMPILEKERNELS + +#define GPUCA_RTC_SPECIAL_CODE(...) GPUCA_RTC_SPECIAL_CODE(__VA_ARGS__) +#define GPUCA_DETERMINISTIC_CODE(...) GPUCA_DETERMINISTIC_CODE(__VA_ARGS__) +#define GPUCA_RTC_CONSTEXPR GPUCA_RTC_CONSTEXPR + +#include "GPUReconstructionIncludesDeviceAll.h" + +#ifndef GPUCA_GPUCODE_DEVICE +#error RTC Preprocessing must run on device code +#endif diff --git a/GPU/GPUTracking/CMakeLists.txt b/GPU/GPUTracking/CMakeLists.txt index cc44b2003a01d..0467b060ead4f 100644 --- a/GPU/GPUTracking/CMakeLists.txt +++ b/GPU/GPUTracking/CMakeLists.txt @@ -434,6 +434,10 @@ install(DIRECTORY utils set(GPUCA_COMPILER_VERSIONS "CUDA ${CMAKE_CUDA_COMPILER_VERSION} HIP ${hip_VERSION} ${CMAKE_HIP_COMPILER_VERSION} ROCM ${CMAKE_HIP_COMPILER_ROCM_ROOT} LLVM ${LLVM_PACKAGE_VERSION}") +if(ENABLE_CPUDEVICE) + set(CPUDEVICE_ENABLED ON) +endif() + # GPUReconstructionLibrary needs to know which GPU backends are enabled for proper error messages configure_file(Base/GPUReconstructionAvailableBackends.template.h ${CMAKE_CURRENT_BINARY_DIR}/GPUReconstructionAvailableBackends.h) set_source_files_properties(Base/GPUReconstructionLibrary.cxx @@ -498,6 +502,9 @@ elseif(CUDA_ENABLED OR OPENCL_ENABLED OR HIP_ENABLED) add_subdirectory(Base/hip) endif() endif() +if(CPUDEVICE_ENABLED) + add_subdirectory(Base/cpudevice) +endif() if(ALIGPU_BUILD_TYPE STREQUAL "O2" OR ALIGPU_BUILD_TYPE STREQUAL "Standalone") if(NOT ${CMAKE_SYSTEM} MATCHES Darwin) diff --git a/GPU/GPUTracking/DataTypes/GPUDataTypesConfig.h b/GPU/GPUTracking/DataTypes/GPUDataTypesConfig.h index 8fc77dec6dbe2..82302605adb32 100644 --- a/GPU/GPUTracking/DataTypes/GPUDataTypesConfig.h +++ b/GPU/GPUTracking/DataTypes/GPUDataTypesConfig.h @@ -32,7 +32,7 @@ namespace gpudatatypes { // clang-format off enum class GeometryType : uint32_t { RESERVED_GEOMETRY = 0, ALIROOT = 1, O2 = 2 }; -enum DeviceType : uint32_t { INVALID_DEVICE = 0, CPU = 1, CUDA = 2, HIP = 3, OCL = 4, METAL = 5 }; +enum DeviceType : uint32_t { INVALID_DEVICE = 0, CPU = 1, CUDA = 2, HIP = 3, OCL = 4, METAL = 5, CPUDEVICE = 6 }; enum class GeneralStep : uint32_t { Prepare = 1, QA = 2 }; // clang-format on @@ -58,7 +58,7 @@ enum class InOutType : uint32_t { TPCClusters = 1, ITSClusters = 128, ITSTracks = 256 }; #if !defined(__OPENCL__) && !defined(__METAL__) -static constexpr const char* const DEVICE_TYPE_NAMES[] = {"INVALID", "CPU", "CUDA", "HIP", "OCL", "METAL"}; +static constexpr const char* const DEVICE_TYPE_NAMES[] = {"INVALID", "CPU", "CUDA", "HIP", "OCL", "METAL", "CPUDEVICE"}; static constexpr const char* const RECO_STEP_NAMES[] = {"TPC Transformation", "TPC Sector Tracking", "TPC Track Merging and Fit", "TPC Compression", "TRD Tracking", "ITS Tracking", "TPC dEdx Computation", "TPC Cluster Finding", "TPC Decompression", "Global Refit"}; static constexpr const char* const GENERAL_STEP_NAMES[] = {"Prepare", "QA"}; constexpr static int32_t N_RECO_STEPS = sizeof(gpudatatypes::RECO_STEP_NAMES) / sizeof(gpudatatypes::RECO_STEP_NAMES[0]); diff --git a/GPU/GPUTracking/Definitions/GPUSettingsList.h b/GPU/GPUTracking/Definitions/GPUSettingsList.h index 3c479646bf844..b82eee36c43a5 100644 --- a/GPU/GPUTracking/Definitions/GPUSettingsList.h +++ b/GPU/GPUTracking/Definitions/GPUSettingsList.h @@ -605,7 +605,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 / OCL) or CPU or AUTO") +AddOption(gpuType, std::string, "AUTO", "", 0, "GPU type (CUDA / HIP / OCL / CPUDEVICE) or CPU or AUTO") AddOption(runGPUforce, bool, true, "", 0, "Force usage of the specified GPU device type, no CPU fallback") AddOption(noprompt, bool, true, "", 0, "Do prompt for keypress before exiting") AddOption(continueOnError, bool, false, "", 0, "Continue processing after an error") @@ -672,7 +672,7 @@ AddOption(constBz, bool, false, "", 0, "force constant Bz for tests") 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(overrideNHbfPerTF, int32_t, 0, "", 0, "Overrides the number of HBF per TF if != 0") AddOption(overrideTPCTimeBinCur, int32_t, 0, "", 0, "Overrides TPC time bin cut if > 0") -AddOption(deviceType, std::string, "CPU", "", 0, "Device type, CPU | CUDA | HIP | OCL | METAL") +AddOption(deviceType, std::string, "CPU", "", 0, "Device type, CPU | CUDA | HIP | OCL | METAL | CPUDEVICE") AddOption(forceDeviceType, bool, true, "", 0, "force device type, otherwise allows fall-back to CPU") AddOption(synchronousProcessing, bool, false, "", 0, "Apply performance shortcuts for synchronous processing, disable unneeded steps") AddOption(dump, int32_t, 0, "", 0, "Dump events for standalone benchmark: 1 = dump events, 2 = dump events and skip processing in workflow") diff --git a/GPU/GPUTracking/Definitions/clusterFinderDefs.h b/GPU/GPUTracking/Definitions/clusterFinderDefs.h index e0b080b1abc08..dea7b146f9344 100644 --- a/GPU/GPUTracking/Definitions/clusterFinderDefs.h +++ b/GPU/GPUTracking/Definitions/clusterFinderDefs.h @@ -23,7 +23,7 @@ #define SCRATCH_PAD_SEARCH_N 8 // TODO: Change all this to constexpr where possible, like in GPUDefParametersConstants.h #define SCRATCH_PAD_COUNT_N 16 -#if defined(GPUCA_GPUCODE) +#if defined(GPUCA_GPUCODE) && !defined(GPUCA_DEVICE_SINGLE_THREAD) #define SCRATCH_PAD_BUILD_N 8 #define SCRATCH_PAD_NOISE_N 8 #else diff --git a/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx b/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx index 2d3449ca0877b..1aa9690453b87 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx @@ -554,7 +554,7 @@ void GPUChainTracking::RunTPCClusterizer_compactPeaks(GPUTPCClusterFinder& clust { auto& in = stage ? clustererShadow.mPpeakPositions : clustererShadow.mPpositions; auto& out = stage ? clustererShadow.mPfilteredPeakPositions : clustererShadow.mPpeakPositions; - if (doGPU) { + if (doGPU && mRec->getGPUParameters(doGPU).par_CF_SCAN_WORKGROUP_SIZE > 1) { // In case we have only 1 device thread, run the CPU version const uint32_t iSector = clusterer.mISector; auto& count = stage ? clusterer.mPmemory->counters.nPeaks : clusterer.mPmemory->counters.nPositions; @@ -589,11 +589,14 @@ void GPUChainTracking::RunTPCClusterizer_compactPeaks(GPUTPCClusterFinder& clust runKernel({GetGrid(count, scanWorkgroupSize, lane), {iSector}}, 1, stage, in, out); } else { - auto& nOut = stage ? clusterer.mPmemory->counters.nClusters : clusterer.mPmemory->counters.nPeaks; - auto& nIn = stage ? clusterer.mPmemory->counters.nPeaks : clusterer.mPmemory->counters.nPositions; + if (doGPU) { + SynchronizeStream(lane); // doGPU with CF_SCAN_WORKGROUP_SIZE <= 1 + } + auto& nOut = stage ? clustererShadow.mPmemory->counters.nClusters : clustererShadow.mPmemory->counters.nPeaks; + auto& nIn = stage ? clustererShadow.mPmemory->counters.nPeaks : clustererShadow.mPmemory->counters.nPositions; size_t count = 0; for (size_t i = 0; i < nIn; i++) { - if (clusterer.mPisPeak[i]) { + if (clustererShadow.mPisPeak[i]) { out[count++] = in[i]; } } diff --git a/GPU/GPUTracking/Standalone/cmake/config.cmake b/GPU/GPUTracking/Standalone/cmake/config.cmake index 5a899945dac0a..8a008245745cb 100644 --- a/GPU/GPUTracking/Standalone/cmake/config.cmake +++ b/GPU/GPUTracking/Standalone/cmake/config.cmake @@ -12,9 +12,10 @@ # This is the configuration file for the standalone build # Its options do not affect the O2 build !!!! -set(ENABLE_CUDA AUTO) -set(ENABLE_HIP AUTO) -set(ENABLE_OPENCL AUTO) +set(ENABLE_CUDA AUTO) # Enable CUDA compilation. Possible Settings are 1/N, 0/OFF, AUTO (for auto-detection) +set(ENABLE_HIP AUTO) # Enable AMD ROCm / HIP compilation (same options as CUDA) +set(ENABLE_OPENCL AUTO) # Enable OpenCL compilation (same options as CUDA) +set(ENABLE_CPUDEVICE 1) # Enable CPUDevice backend: runs the CPU kernels on the host, but behaves like a GPU device (for debugging) set(GPUCA_CONFIG_VC 1) set(GPUCA_CONFIG_FMT 1) set(GPUCA_CONFIG_ROOT 1) diff --git a/GPU/GPUTracking/TPCClusterFinder/CfUtils.h b/GPU/GPUTracking/TPCClusterFinder/CfUtils.h index f9e3f7a304d84..3bc426522deff 100644 --- a/GPU/GPUTracking/TPCClusterFinder/CfUtils.h +++ b/GPU/GPUTracking/TPCClusterFinder/CfUtils.h @@ -177,7 +177,7 @@ class CfUtils const CfChargePos* posBcast, GPUgeneric() T* buf) { -#if defined(GPUCA_GPUCODE) +#if defined(GPUCA_GPUCODE) && !defined(GPUCA_DEVICE_SINGLE_THREAD) // Cooperative load by the threads of the block GPUbarrier(); uint16_t x = ll % N; uint16_t y = ll / N; @@ -222,7 +222,7 @@ class CfUtils const uint8_t* aboveThreshold, GPUgeneric() T* buf) { -#if defined(GPUCA_GPUCODE) +#if defined(GPUCA_GPUCODE) && !defined(GPUCA_DEVICE_SINGLE_THREAD) // Cooperative load by the threads of the block GPUbarrier(); uint16_t y = ll / N; uint16_t x = ll % N; diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx index 14349682a6fba..ecd11250b239c 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx @@ -230,7 +230,7 @@ GPUd() void GPUTPCCFCheckPadBaseline::Thread<0>(int32_t nBlocks, int32_t nThread // Blocks require `nextMultipleOf<64>(138 * 4) = 576` threads to process the largest TPC rows with 138 pads correctly. GPUd() void GPUTPCCFCheckPadBaseline::CheckBaselineGPU(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer) { -#ifdef GPUCA_GPUCODE +#if defined(GPUCA_GPUCODE) && !defined(GPUCA_DEVICE_SINGLE_THREAD) // Needs 576 threads per block, CPUDevice RTC uses the host-compiled kernel static_assert(GPUCA_GET_THREAD_COUNT(GPUCA_LB_GPUTPCCFCheckPadBaseline) == 576); if (iBlock >= (int32_t)GPUTPCGeometry::NROWS) { return; diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc index cf6a065743a22..b1c1aef196e37 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc @@ -208,7 +208,7 @@ GPUdii() void GPUTPCCFClusterizer::buildCluster( labelAcc); } -#if defined(GPUCA_GPUCODE) +#if defined(GPUCA_GPUCODE) && !defined(GPUCA_DEVICE_SINGLE_THREAD) CfUtils::condBlockLoad( chargeMap, wgSizeHalf, diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.cxx index 1237e1fe5930f..ea7e0c577031f 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.cxx @@ -263,7 +263,7 @@ GPUd() void GPUTPCCFNoiseSuppression::findMinimaAndPeaks( bigger); } -#if defined(GPUCA_GPUCODE) +#if defined(GPUCA_GPUCODE) && !defined(GPUCA_DEVICE_SINGLE_THREAD) CfUtils::blockLoad( chargeMap, wgSizeHalf,