Skip to content

Commit a5caa27

Browse files
davidrohrdavidrohr
committed
GPU: Count kernel number in Cmake, and pass to MACROS, to simplify the preprocessor logic
1 parent a0e63ef commit a5caa27

File tree

7 files changed

+24
-38
lines changed

7 files changed

+24
-38
lines changed

‎GPU/GPUTracking/Base/GPUReconstructionCPU.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -138,7 +138,7 @@ gpu_reconstruction_kernels::krnlProperties GPUReconstructionCPU::getKernelProper
138138
return ret;
139139
}
140140

141-
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types) \
141+
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types, ...) \
142142
template void GPUReconstructionCPUBackend::runKernelBackend<GPUCA_M_KRNL_TEMPLATE(x_class)>(const krnlSetupArgs<GPUCA_M_KRNL_TEMPLATE(x_class) GPUCA_M_STRIP(x_types)>& args); \
143143
template krnlProperties GPUReconstructionCPU::getKernelProperties<GPUCA_M_KRNL_TEMPLATE(x_class)>(int gpu);
144144
#include "GPUReconstructionKernelList.h"

‎GPU/GPUTracking/Base/GPUReconstructionCPU.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -72,7 +72,7 @@ class GPUReconstructionCPU : public GPUReconstructionKernels<GPUReconstructionCP
7272

7373
GPUReconstructionCPU(const GPUSettingsDeviceBackend& cfg) : GPUReconstructionKernels(cfg) {}
7474

75-
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types) \
75+
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types, ...) \
7676
inline void runKernelImplWrapper(gpu_reconstruction_kernels::classArgument<GPUCA_M_KRNL_TEMPLATE(x_class)>, bool cpuFallback, double& timer, krnlSetup&& setup GPUCA_M_STRIP(x_arguments)) \
7777
{ \
7878
if (cpuFallback) { \

‎GPU/GPUTracking/Base/GPUReconstructionProcessing.cxx‎

Lines changed: 14 additions & 30 deletions
Original file line numberDiff line numberDiff line change
@@ -143,38 +143,22 @@ std::unique_ptr<gpu_reconstruction_kernels::threadContext> GPUReconstructionProc
143143
gpu_reconstruction_kernels::threadContext::threadContext() = default;
144144
gpu_reconstruction_kernels::threadContext::~threadContext() = default;
145145

146-
template <class T, int32_t I>
147-
uint32_t GPUReconstructionProcessing::GetKernelNum(int32_t k)
148-
{
149-
static int32_t num = k;
150-
if (num < 0) {
151-
throw std::runtime_error("Internal Error - Kernel Number not Set");
152-
}
153-
return num;
154-
}
155-
156-
namespace o2::gpu::internal
157-
{
158-
static std::vector<std::string> initKernelNames()
159-
{
160-
std::vector<std::string> retVal;
161-
#define GPUCA_KRNL(x_class, ...) \
162-
GPUReconstructionProcessing::GetKernelNum<GPUCA_M_KRNL_TEMPLATE(x_class)>(retVal.size()); \
163-
retVal.emplace_back(GPUCA_M_STR(GPUCA_M_KRNL_NAME(x_class)));
146+
const std::vector<std::string> GPUReconstructionProcessing::mKernelNames = {
147+
#define GPUCA_KRNL(x_class, ...) GPUCA_M_STR(GPUCA_M_KRNL_NAME(x_class)),
164148
#include "GPUReconstructionKernelList.h"
165149
#undef GPUCA_KRNL
166-
return retVal;
167-
}
168-
} // namespace o2::gpu::internal
169-
170-
const std::vector<std::string> GPUReconstructionProcessing::mKernelNames = o2::gpu::internal::initKernelNames();
171-
172-
#define GPUCA_KRNL(x_class, ...) \
173-
template uint32_t GPUReconstructionProcessing::GetKernelNum<GPUCA_M_KRNL_TEMPLATE(x_class)>(int32_t); \
174-
template <> \
175-
const char* GPUReconstructionProcessing::GetKernelName<GPUCA_M_KRNL_TEMPLATE(x_class)>() \
176-
{ \
177-
return GPUCA_M_STR(GPUCA_M_KRNL_NAME(x_class)); \
150+
};
151+
152+
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types, x_num) \
153+
template <> \
154+
uint32_t GPUReconstructionProcessing::GetKernelNum<GPUCA_M_KRNL_TEMPLATE(x_class)>() \
155+
{ \
156+
return x_num; \
157+
} \
158+
template <> \
159+
const char* GPUReconstructionProcessing::GetKernelName<GPUCA_M_KRNL_TEMPLATE(x_class)>() \
160+
{ \
161+
return GPUCA_M_STR(GPUCA_M_KRNL_NAME(x_class)); \
178162
}
179163
#include "GPUReconstructionKernelList.h"
180164
#undef GPUCA_KRNL

‎GPU/GPUTracking/Base/GPUReconstructionProcessing.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -80,7 +80,7 @@ class GPUReconstructionProcessing : public GPUReconstruction
8080
static const char* GetKernelName();
8181
const std::string& GetKernelName(int32_t i) const { return mKernelNames[i]; }
8282
template <class T, int32_t I = 0>
83-
static uint32_t GetKernelNum(int32_t k = -1);
83+
static uint32_t GetKernelNum();
8484

8585
// Public queries for timers
8686
auto& getRecoStepTimer(RecoStep step) { return mTimersRecoSteps[getRecoStepNum(step)]; }

‎GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAKernels.cu‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -83,14 +83,14 @@ void GPUReconstructionCUDABackend::runKernelBackend(const krnlSetupArgs<T, I, Ar
8383
#define GPUCA_KRNL_REG(args) __launch_bounds__(GPUCA_M_MAX2_3(GPUCA_M_STRIP(args)))
8484

8585
#if defined(GPUCA_KERNEL_COMPILE_MODE) && GPUCA_KERNEL_COMPILE_MODE == 1 // ---------- COMPILE_MODE = perkernel ----------
86-
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types) template void GPUReconstructionCUDABackend::runKernelBackend<GPUCA_M_KRNL_TEMPLATE(x_class)>(const krnlSetupArgs<GPUCA_M_KRNL_TEMPLATE(x_class) GPUCA_M_STRIP(x_types)>& args);
86+
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types, ...) template void GPUReconstructionCUDABackend::runKernelBackend<GPUCA_M_KRNL_TEMPLATE(x_class)>(const krnlSetupArgs<GPUCA_M_KRNL_TEMPLATE(x_class) GPUCA_M_STRIP(x_types)>& args);
8787
#else // ---------- COMPILE_MODE = onefile | rdc ----------
8888
#if defined(GPUCA_KERNEL_COMPILE_MODE) && GPUCA_KERNEL_COMPILE_MODE == 2
8989
#define GPUCA_KRNL_DEFONLY // COMPILE_MODE = rdc
9090
#endif
9191

92-
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types) \
93-
GPUCA_KRNL_HOST(x_class, x_attributes, x_arguments, x_forward, x_types) \
92+
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types, ...) \
93+
GPUCA_KRNL_HOST(x_class, x_attributes, x_arguments, x_forward, x_types, __VA_ARGS__) \
9494
template void GPUReconstructionCUDABackend::runKernelBackend<GPUCA_M_KRNL_TEMPLATE(x_class)>(const krnlSetupArgs<GPUCA_M_KRNL_TEMPLATE(x_class) GPUCA_M_STRIP(x_types)>& args);
9595

9696
#ifndef __HIPCC__ // CUDA version

‎GPU/GPUTracking/Base/opencl/GPUReconstructionOCLKernels.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -91,6 +91,6 @@ int32_t GPUReconstructionOCLBackend::AddKernels()
9191
return 0;
9292
}
9393

94-
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types) template void GPUReconstructionOCLBackend::runKernelBackend<GPUCA_M_KRNL_TEMPLATE(x_class)>(const krnlSetupArgs<GPUCA_M_KRNL_TEMPLATE(x_class) GPUCA_M_STRIP(x_types)>& args);
94+
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types, ...) template void GPUReconstructionOCLBackend::runKernelBackend<GPUCA_M_KRNL_TEMPLATE(x_class)>(const krnlSetupArgs<GPUCA_M_KRNL_TEMPLATE(x_class) GPUCA_M_STRIP(x_types)>& args);
9595
#include "GPUReconstructionKernelList.h"
9696
#undef GPUCA_KRNL

‎GPU/GPUTracking/cmake/kernel_helpers.cmake‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -64,7 +64,9 @@ function(o2_gpu_add_kernel kernel_name kernel_files)
6464
endif()
6565
set(TMP_PRE "")
6666
set(TMP_POST "")
67-
set(TMP_KERNEL "GPUCA_KRNL${TMP_BOUNDS}((${kernel_name}), (${kernel_extra}), (${OPT1}), (${OPT2}), (${OPT3}))\n")
67+
get_property(LIST_KERNELS TARGET O2_GPU_KERNELS PROPERTY O2_GPU_KERNELS)
68+
list(LENGTH LIST_KERNELS KERNEL_COUNT)
69+
set(TMP_KERNEL "GPUCA_KRNL${TMP_BOUNDS}((${kernel_name}), (${kernel_extra}), (${OPT1}), (${OPT2}), (${OPT3}), ${KERNEL_COUNT})\n")
6870
separate_arguments(kernel_files NATIVE_COMMAND ${kernel_files})
6971
list(GET kernel_files 0 TMP_KERNEL_CLASS_FILE)
7072
if (TMP_KERNEL_CLASS_FILE STREQUAL "=")

0 commit comments

Comments
 (0)