Skip to content

Commit c44cd25

Browse files
committed
GPU: Add CPUDevice backend to emulate device-like processing on host
1 parent da0be42 commit c44cd25

25 files changed

Lines changed: 1189 additions & 25 deletions

‎GPU/Common/GPUCommonAlgorithm.h‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -17,8 +17,10 @@
1717

1818
#include "GPUCommonDef.h"
1919

20-
#if !defined(GPUCA_GPUCODE) // Could also enable custom search on the CPU, but it is not always faster, so we stick to std::sort
20+
#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
21+
#ifndef GPUCA_GPUCODE_DEVICE
2122
#include <algorithm>
23+
#endif
2224
#define GPUCA_ALGORITHM_STD
2325
#endif
2426

@@ -280,7 +282,7 @@ GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end GPUCA_THREAD_INFO_
280282
template <class T, class S>
281283
GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end, const S& comp GPUCA_THREAD_INFO_DECL)
282284
{
283-
#ifndef GPUCA_GPUCODE
285+
#if !defined(GPUCA_GPUCODE) || defined(GPUCA_DEVICE_SINGLE_THREAD)
284286
GPUCommonAlgorithm::sort(begin, end, comp);
285287
#elif defined(GPUCA_DETERMINISTIC_MODE) // Not using GPUCA_DETERMINISTIC_CODE, which is enforced in TPC compression
286288
if (get_local_id(0) == 0) {

‎GPU/Common/GPUCommonArray.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -15,7 +15,7 @@
1515
#ifndef GPUCOMMONARRAY_H
1616
#define GPUCOMMONARRAY_H
1717

18-
#if !defined(GPUCA_GPUCODE_DEVICE) || defined(__CUDACC__) || defined(__HIPCC__) // TODO: Get rid of GPUCommonArray once OpenCL supports <array>
18+
#if !defined(GPUCA_GPUCODE_DEVICE) || defined(__CUDACC__) || defined(__HIPCC__) || defined(__CPUDEVICE__) // TODO: Get rid of GPUCommonArray once OpenCL supports <array>
1919
#ifndef GPUCA_GPUCODE_COMPILEKERNELS
2020
#include <array>
2121
#endif

‎GPU/Common/GPUCommonDef.h‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -31,12 +31,12 @@
3131
#include "GPUCommonDefSettings.h"
3232

3333
#if !defined(__CLING__) && !defined(G__ROOT) // No GPU code for ROOT
34-
#if defined(__CUDACC__) || defined(__OPENCL__) || defined(__HIPCC__) || defined(__OPENCL_HOST__) || defined(__METAL__) || defined(__METAL_HOST__)
34+
#if defined(__CUDACC__) || defined(__OPENCL__) || defined(__HIPCC__) || defined(__OPENCL_HOST__) || defined(__METAL__) || defined(__METAL_HOST__) || defined(__CPUDEVICE__)
3535
#define GPUCA_GPUCODE // Compiled by GPU compiler
3636
#endif
3737

3838
#if defined(GPUCA_GPUCODE)
39-
#if defined(__CUDA_ARCH__) || defined(__OPENCL__) || defined(__HIP_DEVICE_COMPILE__) || defined(__METAL_VERSION__)
39+
#if defined(__CUDA_ARCH__) || defined(__OPENCL__) || defined(__HIP_DEVICE_COMPILE__) || defined(__METAL_VERSION__) || defined(__CPUDEVICE__)
4040
#define GPUCA_GPUCODE_DEVICE // Executed on device
4141
#endif
4242
#if defined(__CUDACC__)
@@ -47,6 +47,8 @@
4747
#define GPUCA_GPUTYPE OCL
4848
#elif defined(__METAL__) || defined(__METAL_HOST__)
4949
#define GPUCA_GPUTYPE METAL
50+
#elif defined(__CPUDEVICE__)
51+
#define GPUCA_GPUTYPE CPUDEVICE
5052
#endif
5153
#endif
5254
#endif

‎GPU/Common/GPUCommonDefAPI.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,7 @@
2727
//Define macros for GPU keywords. i-version defines inline functions.
2828
//All host-functions in GPU code are automatically inlined, to avoid duplicate symbols.
2929
//For non-inline host only functions, use no keyword at all!
30-
#if !defined(GPUCA_GPUCODE) || defined(__OPENCL_HOST__) || defined(__METAL_HOST__) // For host / ROOT dictionary
30+
#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
3131
#define GPUd() // device function
3232
#define GPUdDefault() // default (constructor / operator) device function
3333
#define GPUhdDefault() // default (constructor / operator) host device function

‎GPU/Common/GPUCommonMath.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -236,7 +236,7 @@ GPUdi() constexpr T GPUCommonMath::nextMultipleOf(T val)
236236

237237
GPUdi() float2 GPUCommonMath::MakeFloat2(float x, float y)
238238
{
239-
#if !defined(GPUCA_GPUCODE) || defined(__OPENCL__) || defined(__OPENCL_HOST__) || defined(__METAL__) || defined(__METAL_HOST__)
239+
#if !defined(GPUCA_GPUCODE) || defined(__OPENCL__) || defined(__OPENCL_HOST__) || defined(__METAL__) || defined(__METAL_HOST__) || defined(__CPUDEVICE__)
240240
float2 ret = {x, y};
241241
return ret;
242242
#else

‎GPU/GPUTracking/Base/GPUReconstruction.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -427,7 +427,7 @@ class GPUReconstruction
427427
void* mGPULib;
428428
void* mGPUEntry;
429429
};
430-
static std::shared_ptr<LibraryLoader> sLibCUDA, sLibHIP, sLibOCL, sLibMETAL;
430+
static std::shared_ptr<LibraryLoader> sLibCUDA, sLibHIP, sLibOCL, sLibMETAL, sLibCPUDEVICE;
431431

432432
// Debugging
433433
struct debugInternal;

‎GPU/GPUTracking/Base/GPUReconstructionAvailableBackends.template.h‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -17,5 +17,6 @@
1717
#cmakedefine HIP_ENABLED
1818
#cmakedefine OPENCL_ENABLED
1919
#cmakedefine METAL_ENABLED
20+
#cmakedefine CPUDEVICE_ENABLED
2021
#cmakedefine GPUCA_COMPILER_VERSIONS @GPUCA_COMPILER_VERSIONS@
2122
// clang-format on

‎GPU/GPUTracking/Base/GPUReconstructionLibrary.cxx‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -106,6 +106,10 @@ std::shared_ptr<GPUReconstruction::LibraryLoader>* GPUReconstruction::GetLibrary
106106
} else if (type == DeviceType::METAL) {
107107
#ifdef METAL_ENABLED
108108
return &sLibMETAL;
109+
#endif
110+
} else if (type == DeviceType::CPUDEVICE) {
111+
#ifdef CPUDEVICE_ENABLED
112+
return &sLibCPUDEVICE;
109113
#endif
110114
} else {
111115
GPUError("Error: Invalid device type %u", (uint32_t)type);
@@ -131,6 +135,7 @@ std::shared_ptr<GPUReconstruction::LibraryLoader> GPUReconstruction::sLibCUDA(ne
131135
std::shared_ptr<GPUReconstruction::LibraryLoader> GPUReconstruction::sLibHIP(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingHIP" LIBRARY_EXTENSION, "GPUReconstruction_Create_HIP"));
132136
std::shared_ptr<GPUReconstruction::LibraryLoader> GPUReconstruction::sLibOCL(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingOCL" LIBRARY_EXTENSION, "GPUReconstruction_Create_OCL"));
133137
std::shared_ptr<GPUReconstruction::LibraryLoader> GPUReconstruction::sLibMETAL(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingMETAL" LIBRARY_EXTENSION, "GPUReconstruction_Create_METAL"));
138+
std::shared_ptr<GPUReconstruction::LibraryLoader> GPUReconstruction::sLibCPUDEVICE(new GPUReconstruction::LibraryLoader("lib" LIBRARY_PREFIX "GPUTrackingCPUDevice" LIBRARY_EXTENSION, "GPUReconstruction_Create_CPUDEVICE"));
134139

135140
GPUReconstruction::LibraryLoader::LibraryLoader(const char* lib, const char* func) : mLibName(lib), mFuncName(func), mGPULib(nullptr), mGPUEntry(nullptr) {}
136141

Lines changed: 105 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,105 @@
1+
# Copyright 2019-2026 CERN and copyright holders of ALICE O2.
2+
# See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
3+
# All rights not expressly granted are reserved.
4+
#
5+
# This software is distributed under the terms of the GNU General Public
6+
# License v3 (GPL Version 3), copied verbatim in the file "COPYING".
7+
#
8+
# In applying this license CERN does not waive the privileges and immunities
9+
# granted to it by virtue of its status as an Intergovernmental Organization
10+
# or submit itself to any jurisdiction.
11+
12+
set(MODULE GPUTrackingCPUDevice)
13+
14+
set(SRCS GPUReconstructionCPUDevice.cxx GPUReconstructionCPUDeviceGenRTC.cxx)
15+
set(HDRS GPUReconstructionCPUDevice.h GPUReconstructionCPUDeviceInternals.h GPUReconstructionCPUDeviceIncludesSystem.h)
16+
17+
if(CMAKE_SYSTEM_NAME STREQUAL "Linux")
18+
set(GPUCA_CPUDEVICE_RTC ON)
19+
else()
20+
set(GPUCA_CPUDEVICE_RTC OFF)
21+
endif()
22+
message(STATUS "Building GPUTracking with CPUDevice support (RTC: ${GPUCA_CPUDEVICE_RTC})")
23+
24+
if(ALIGPU_BUILD_TYPE STREQUAL "O2")
25+
set(defineIncludeSrc "O2::${MODULE}")
26+
else()
27+
set(defineIncludeSrc "${MODULE}")
28+
endif()
29+
30+
if(GPUCA_CPUDEVICE_RTC) # RTC compiles the kernels with the host compiler and the host compiler flags, and loads them via dlopen
31+
enable_language(ASM)
32+
if(ALIGPU_BUILD_TYPE STREQUAL "Standalone")
33+
get_filename_component(GPUDIR ${CMAKE_SOURCE_DIR}/../ ABSOLUTE)
34+
else()
35+
set(GPUDIR ${CMAKE_SOURCE_DIR}/GPU/GPUTracking)
36+
endif()
37+
string(TOUPPER "${CMAKE_BUILD_TYPE}" TMP_BUILD_TYPE_UPPER)
38+
set(GPU_RTC_DEFINES "-D$<JOIN:$<TARGET_PROPERTY:${defineIncludeSrc},COMPILE_DEFINITIONS>,$<SEMICOLON>-D>")
39+
set(GPU_RTC_INCLUDES "-I$<JOIN:$<FILTER:$<TARGET_PROPERTY:${defineIncludeSrc},INCLUDE_DIRECTORIES>,EXCLUDE,^/usr/include/?>,$<SEMICOLON>-I>"
40+
-I${CMAKE_SOURCE_DIR}/Detectors/Base/src
41+
-I${CMAKE_SOURCE_DIR}/Detectors/TRD/base/src
42+
)
43+
separate_arguments(GPU_RTC_FLAGS NATIVE_COMMAND "-std=c++${CMAKE_CXX_STANDARD} ${CMAKE_CXX_FLAGS} ${CMAKE_CXX_FLAGS_${TMP_BUILD_TYPE_UPPER}} -w")
44+
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)
45+
set(GPU_RTC_SRC ${GPUDIR}/Base/cpudevice/GPUReconstructionCPUDevicertc.cxx)
46+
set(GPU_RTC_BIN ${CMAKE_CURRENT_BINARY_DIR}/GPUReconstructionCPUDevicertc)
47+
48+
add_custom_command(
49+
OUTPUT ${GPU_RTC_BIN}.src
50+
COMMAND cp ${GPUDIR}/Base/GPUStdSystemHeaders.h ${GPU_RTC_BIN}.src
51+
COMMAND cat ${GPUDIR}/Base/cpudevice/GPUReconstructionCPUDeviceIncludesSystem.h | grep -v GPUStdSystemHeaders.h >> ${GPU_RTC_BIN}.src
52+
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
53+
DEPENDS ${GPU_RTC_SRC} ${GPUDIR}/Base/GPUStdSystemHeaders.h ${GPUDIR}/Base/cpudevice/GPUReconstructionCPUDeviceIncludesSystem.h
54+
DEPFILE ${GPU_RTC_BIN}.src.d
55+
COMMAND_EXPAND_LISTS
56+
COMMENT "Preparing CPUDevice RTC source file ${GPU_RTC_BIN}.src"
57+
)
58+
add_custom_command(
59+
OUTPUT ${GPU_RTC_BIN}.command
60+
COMMAND echo -n "${CMAKE_CXX_COMPILER} ${GPU_RTC_FLAGS} ${GPU_RTC_LINK_FLAGS}" > ${GPU_RTC_BIN}.command
61+
COMMAND_EXPAND_LISTS VERBATIM
62+
COMMENT "Preparing CPUDevice RTC command file ${GPU_RTC_BIN}.command"
63+
)
64+
add_custom_command(
65+
OUTPUT ${GPU_RTC_BIN}.command.no_fast_math
66+
COMMAND echo -n "${GPUCA_CXX_NO_FAST_MATH_FLAGS}" > ${GPU_RTC_BIN}.command.no_fast_math
67+
COMMAND_EXPAND_LISTS VERBATIM
68+
COMMENT "Preparing CPUDevice RTC NO_FAST_MATH command file ${GPU_RTC_BIN}.command.no_fast_math"
69+
)
70+
foreach(TMP_RESOURCE src command command.no_fast_math)
71+
create_binary_resource(${GPU_RTC_BIN}.${TMP_RESOURCE} ${GPU_RTC_BIN}.${TMP_RESOURCE}.o)
72+
list(APPEND SRCS ${GPU_RTC_BIN}.${TMP_RESOURCE}.o)
73+
endforeach()
74+
endif()
75+
76+
if(ALIGPU_BUILD_TYPE STREQUAL "O2")
77+
o2_add_library(
78+
${MODULE}
79+
SOURCES ${SRCS}
80+
PUBLIC_LINK_LIBRARIES O2::GPUTracking
81+
PRIVATE_INCLUDE_DIRECTORIES
82+
${CMAKE_SOURCE_DIR}/Detectors/Base/src
83+
${CMAKE_SOURCE_DIR}/Detectors/TRD/base/src
84+
${CMAKE_SOURCE_DIR}/DataFormats/Reconstruction/src
85+
${CMAKE_CURRENT_SOURCE_DIR}
86+
TARGETVARNAME targetName)
87+
endif()
88+
89+
if(ALIGPU_BUILD_TYPE STREQUAL "Standalone")
90+
set(targetName "${MODULE}")
91+
add_library(${MODULE} SHARED ${SRCS})
92+
add_library(O2::${MODULE} ALIAS ${MODULE})
93+
target_link_libraries(${MODULE} PUBLIC O2::GPUTracking)
94+
install(TARGETS ${MODULE})
95+
include_directories(${CMAKE_CURRENT_SOURCE_DIR})
96+
endif()
97+
98+
install(FILES ${HDRS} DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/GPU)
99+
100+
target_compile_definitions(${targetName} PRIVATE $<TARGET_PROPERTY:O2::GPUTracking,COMPILE_DEFINITIONS>)
101+
target_link_libraries(${targetName} PRIVATE TBB::tbb)
102+
if(GPUCA_CPUDEVICE_RTC)
103+
target_compile_definitions(${targetName} PRIVATE GPUCA_CPUDEVICE_RTC)
104+
target_link_libraries(${targetName} PRIVATE ${CMAKE_DL_LIBS})
105+
endif()

0 commit comments

Comments
 (0)