Skip to content

Commit 52c2328

Browse files
davidrohrdavidrohr
committed
GPU: Get rid of backendInternal additional wrapper
1 parent 29e2526 commit 52c2328

11 files changed

+59
-77
lines changed

‎GPU/GPUTracking/Base/GPUReconstructionCPU.cxx‎

Lines changed: 3 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -54,7 +54,7 @@ GPUReconstructionCPU::~GPUReconstructionCPU()
5454
}
5555

5656
template <class T, int32_t I, typename... Args>
57-
inline void GPUReconstructionCPU::runKernelBackendInternal(const krnlSetupTime& _xyz, const Args&... args)
57+
inline void GPUReconstructionCPU::runKernelBackend(const krnlSetupTime& _xyz, const Args&... args)
5858
{
5959
auto& x = _xyz.x;
6060
auto& y = _xyz.y;
@@ -88,7 +88,7 @@ inline void GPUReconstructionCPU::runKernelBackendInternal(const krnlSetupTime&
8888
}
8989

9090
template <>
91-
inline void GPUReconstructionCPU::runKernelBackendInternal<GPUMemClean16, 0>(const krnlSetupTime& _xyz, void* const& ptr, uint64_t const& size)
91+
inline void GPUReconstructionCPU::runKernelBackend<GPUMemClean16, 0>(const krnlSetupTime& _xyz, void* const& ptr, uint64_t const& size)
9292
{
9393
int32_t nThreads = std::max<int32_t>(1, std::min<int32_t>(size / (16 * 1024 * 1024), getNKernelHostThreads(true)));
9494
if (nThreads > 1) {
@@ -108,17 +108,6 @@ inline void GPUReconstructionCPU::runKernelBackendInternal<GPUMemClean16, 0>(con
108108
}
109109
}
110110

111-
template <class T, int32_t I, typename... Args>
112-
void GPUReconstructionCPU::runKernelBackend(const krnlSetupArgs<T, I, Args...>& args)
113-
{
114-
#pragma GCC diagnostic push
115-
#if defined(__clang__)
116-
#pragma GCC diagnostic ignored "-Wunused-lambda-capture" // this is not alway captured below
117-
#endif
118-
std::apply([this, &args](auto&... vals) { runKernelBackendInternal<T, I, Args...>(args.s, vals...); }, args.v);
119-
#pragma GCC diagnostic push
120-
}
121-
122111
template <class S, int32_t I>
123112
GPUReconstructionProcessing::krnlProperties GPUReconstructionCPU::getKernelProperties(int gpu)
124113
{
@@ -137,8 +126,7 @@ GPUReconstructionProcessing::krnlProperties GPUReconstructionCPU::getKernelPrope
137126
return ret;
138127
}
139128

140-
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types, ...) \
141-
template void GPUReconstructionCPU::runKernelBackend<GPUCA_M_KRNL_TEMPLATE(x_class)>(const krnlSetupArgs<GPUCA_M_KRNL_TEMPLATE(x_class) GPUCA_M_STRIP(x_types)>& args); \
129+
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types, ...) \
142130
template GPUReconstructionProcessing::krnlProperties GPUReconstructionCPU::getKernelProperties<GPUCA_M_KRNL_TEMPLATE(x_class)>(int gpu);
143131
#include "GPUReconstructionKernelList.h"
144132
#undef GPUCA_KRNL

‎GPU/GPUTracking/Base/GPUReconstructionCPU.h‎

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -40,7 +40,7 @@ class GPUReconstructionCPU : public GPUReconstructionProcessing::KernelInterface
4040
template <class S, int32_t I = 0>
4141
krnlProperties getKernelProperties(int gpu = -1);
4242
template <class T, int32_t I = 0, typename... Args>
43-
void runKernelBackend(const krnlSetupArgs<T, I, Args...>& args);
43+
void runKernelBackend(const krnlSetupTime& _xyz, const Args&... args);
4444

4545
virtual int32_t GPUDebug(const char* state = "UNKNOWN", int32_t stream = -1, bool force = false);
4646
int32_t GPUStuck() { return mGPUStuck; }
@@ -59,9 +59,6 @@ class GPUReconstructionCPU : public GPUReconstructionProcessing::KernelInterface
5959

6060
GPUReconstructionCPU(const GPUSettingsDeviceBackend& cfg) : GPUReconstructionProcessing::KernelInterface<GPUReconstructionCPU, GPUReconstructionProcessing>(cfg) {}
6161

62-
template <class T, int32_t I = 0, typename... Args>
63-
void runKernelBackendInternal(const krnlSetupTime& _xyz, const Args&... args);
64-
6562
int32_t registerMemoryForGPU_internal(const void* ptr, size_t size) override { return 0; }
6663
int32_t unregisterMemoryForGPU_internal(const void* ptr) override { return 0; }
6764

‎GPU/GPUTracking/Base/GPUReconstructionProcessingKernels.inc‎

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -21,21 +21,28 @@
2121
namespace o2::gpu
2222
{
2323

24+
#pragma GCC diagnostic push
25+
#if defined(__clang__)
26+
#pragma GCC diagnostic ignored "-Wunused-lambda-capture" // this is not alway captured below
27+
#endif
28+
2429
template <class T, class S>
2530
void GPUReconstructionProcessing::KernelInterface<T, S>::runKernelVirtual(const int num, const void* args)
2631
{
2732
switch (num) { // clang-format off
2833
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types, x_num) \
2934
case x_num: { \
3035
const auto& args2 = *(const krnlSetupArgs<GPUCA_M_KRNL_TEMPLATE(x_class) GPUCA_M_STRIP(x_types)>*)args; \
31-
((T*)this)->template runKernelBackend<GPUCA_M_KRNL_TEMPLATE(x_class)>(args2); \
36+
std::apply([this, &args2](auto&... vals) { ((T*)this)->template runKernelBackend<GPUCA_M_KRNL_TEMPLATE(x_class) GPUCA_M_STRIP(x_types)>(args2.s, vals...); }, args2.v); \
3237
break; \
3338
}
3439
#include "GPUReconstructionKernelList.h"
3540
#undef GPUCA_KRNL
3641
} // clang-format on
3742
}
3843

44+
#pragma GCC diagnostic push
45+
3946
} // namespace o2::gpu
4047

4148
#endif // GPURECONSTRUCTIONPROCESSINGKERNELS_H

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

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,8 @@
2424
#include "GPUParamRTC.h"
2525
#include "GPUReconstructionCUDAHelpers.inc"
2626
#include "GPUDefParametersLoad.inc"
27-
#include "GPUReconstructionProcessingKernels.inc"
27+
#include "GPUReconstructionKernelIncludes.h"
28+
#include "GPUConstantMem.h"
2829

2930
#if defined(GPUCA_KERNEL_COMPILE_MODE) && GPUCA_KERNEL_COMPILE_MODE == 1
3031
#include "utils/qGetLdBinarySymbols.h"

‎GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.h‎

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -45,17 +45,16 @@ class GPUReconstructionCUDA : public GPUReconstructionProcessing::KernelInterfac
4545
virtual int32_t GPUChkErrInternal(const int64_t error, const char* file, int32_t line) const override;
4646

4747
template <class T, int32_t I = 0, typename... Args>
48-
void runKernelBackend(const krnlSetupArgs<T, I, Args...>& args);
48+
void runKernelBackend(const krnlSetupTime& _xyz, const Args&... args);
49+
template <class T, int32_t I = 0, typename... Args>
50+
void runKernelBackendTimed(const krnlSetupTime& _xyz, const Args&... args);
4951

5052
template <class T, class S>
5153
friend GPUh() void GPUCommonAlgorithm::sortOnDevice(auto* rec, int32_t stream, T* begin, size_t N, const S& comp);
5254

5355
protected:
5456
GPUReconstructionCUDAInternals* mInternals;
5557

56-
template <class T, int32_t I = 0, typename... Args>
57-
void runKernelBackendInternal(const krnlSetupTime& _xyz, const Args&... args);
58-
5958
int32_t InitDevice_Runtime() override;
6059
int32_t ExitDevice_Runtime() override;
6160

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

Lines changed: 30 additions & 30 deletions
Original file line numberDiff line numberDiff line change
@@ -23,13 +23,15 @@ using namespace o2::gpu;
2323
#include "GPUReconstructionIncludesDeviceAll.h"
2424

2525
#include "GPUReconstructionCUDAKernelsSpecialize.inc"
26+
#include "GPUReconstructionProcessingKernels.inc"
27+
template void GPUReconstructionProcessing::KernelInterface<GPUReconstructionCUDA, GPUReconstructionDeviceBase>::runKernelVirtual(const int num, const void* args);
2628

2729
#if defined(__HIPCC__) && defined(GPUCA_HAS_GLOBAL_SYMBOL_CONSTANT_MEM)
2830
__global__ void gGPUConstantMemBuffer_dummy(int32_t* p) { *p = *(int32_t*)&gGPUConstantMemBuffer; }
2931
#endif
3032

3133
template <class T, int32_t I, typename... Args>
32-
inline void GPUReconstructionCUDA::runKernelBackendInternal(const krnlSetupTime& _xyz, const Args&... args)
34+
inline void GPUReconstructionCUDA::runKernelBackendTimed(const krnlSetupTime& _xyz, const Args&... args)
3335
{
3436
#if !defined(GPUCA_KERNEL_COMPILE_MODE) || GPUCA_KERNEL_COMPILE_MODE != 1
3537
if (!GetProcessingSettings().rtc.enable) {
@@ -52,18 +54,18 @@ inline void GPUReconstructionCUDA::runKernelBackendInternal(const krnlSetupTime&
5254
}
5355

5456
template <class T, int32_t I, typename... Args>
55-
void GPUReconstructionCUDA::runKernelBackend(const krnlSetupArgs<T, I, Args...>& args)
57+
inline void GPUReconstructionCUDA::runKernelBackend(const krnlSetupTime& _xyz, const Args&... args)
5658
{
57-
auto& x = args.s.x;
58-
auto& z = args.s.z;
59+
auto& x = _xyz.x;
60+
auto& z = _xyz.z;
5961
if (z.evList) {
6062
for (int32_t k = 0; k < z.nEvents; k++) {
6163
GPUChkErr(cudaStreamWaitEvent(mInternals->Streams[x.stream], ((cudaEvent_t*)z.evList)[k], 0));
6264
}
6365
}
6466
{
65-
GPUDebugTiming timer(GetProcessingSettings().deviceTimers && GetProcessingSettings().debugLevel > 0, (deviceEvent*)mDebugEvents, mInternals->Streams, args.s, this);
66-
std::apply([this, &args](auto&... vals) { this->runKernelBackendInternal<T, I, Args...>(args.s, vals...); }, args.v);
67+
GPUDebugTiming timer(GetProcessingSettings().deviceTimers && GetProcessingSettings().debugLevel > 0, (deviceEvent*)mDebugEvents, mInternals->Streams, _xyz, this);
68+
runKernelBackendTimed<T, I, Args...>(_xyz, args...);
6769
}
6870
GPUChkErr(cudaGetLastError());
6971
if (z.ev) {
@@ -74,31 +76,29 @@ void GPUReconstructionCUDA::runKernelBackend(const krnlSetupArgs<T, I, Args...>&
7476
#undef GPUCA_KRNL_REG
7577
#define GPUCA_KRNL_REG(args) __launch_bounds__(GPUCA_M_MAX2_3(GPUCA_M_STRIP(args)))
7678

77-
#if defined(GPUCA_KERNEL_COMPILE_MODE) && GPUCA_KERNEL_COMPILE_MODE == 1 // ---------- COMPILE_MODE = perkernel ----------
78-
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types, ...) template void GPUReconstructionCUDA::runKernelBackend<GPUCA_M_KRNL_TEMPLATE(x_class)>(const krnlSetupArgs<GPUCA_M_KRNL_TEMPLATE(x_class) GPUCA_M_STRIP(x_types)>& args);
79-
#else // ---------- COMPILE_MODE = onefile | rdc ----------
80-
#if defined(GPUCA_KERNEL_COMPILE_MODE) && GPUCA_KERNEL_COMPILE_MODE == 2
81-
#define GPUCA_KRNL_DEFONLY // COMPILE_MODE = rdc
82-
#endif
83-
84-
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types, ...) \
85-
GPUCA_KRNL_HOST(x_class, x_attributes, x_arguments, x_forward, x_types, __VA_ARGS__) \
86-
template void GPUReconstructionCUDA::runKernelBackend<GPUCA_M_KRNL_TEMPLATE(x_class)>(const krnlSetupArgs<GPUCA_M_KRNL_TEMPLATE(x_class) GPUCA_M_STRIP(x_types)>& args);
87-
88-
#ifndef __HIPCC__ // CUDA version
89-
#define GPUCA_KRNL_CALL(x_class, ...) \
90-
GPUCA_M_CAT(krnl_, GPUCA_M_KRNL_NAME(x_class))<<<x.nBlocks, x.nThreads, 0, me->mInternals->Streams[x.stream]>>>(GPUCA_CONSMEM_CALL y.index, args...);
91-
#else // HIP version
92-
#undef GPUCA_KRNL_CUSTOM
93-
#define GPUCA_KRNL_CUSTOM(args) GPUCA_M_STRIP(args)
94-
#define GPUCA_KRNL_CALL(x_class, ...) \
95-
hipLaunchKernelGGL(HIP_KERNEL_NAME(GPUCA_M_CAT(krnl_, GPUCA_M_KRNL_NAME(x_class))), dim3(x.nBlocks), dim3(x.nThreads), 0, me->mInternals->Streams[x.stream], GPUCA_CONSMEM_CALL y.index, args...);
96-
#endif // __HIPCC__
97-
79+
// clang-format off
80+
#if defined(GPUCA_KERNEL_COMPILE_MODE) && GPUCA_KERNEL_COMPILE_MODE != 1 // ---------- COMPILE_MODE = perkernel ----------
81+
#if defined(GPUCA_KERNEL_COMPILE_MODE) && GPUCA_KERNEL_COMPILE_MODE == 2
82+
#define GPUCA_KRNL_DEFONLY // COMPILE_MODE = rdc
83+
#endif
84+
85+
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types, ...) \
86+
GPUCA_KRNL_HOST(x_class, x_attributes, x_arguments, x_forward, x_types, __VA_ARGS__)
87+
88+
#ifndef __HIPCC__ // CUDA version
89+
#define GPUCA_KRNL_CALL(x_class, ...) \
90+
GPUCA_M_CAT(krnl_, GPUCA_M_KRNL_NAME(x_class))<<<x.nBlocks, x.nThreads, 0, me->mInternals->Streams[x.stream]>>>(GPUCA_CONSMEM_CALL y.index, args...);
91+
#else // HIP version
92+
#undef GPUCA_KRNL_CUSTOM
93+
#define GPUCA_KRNL_CUSTOM(args) GPUCA_M_STRIP(args)
94+
#define GPUCA_KRNL_CALL(x_class, ...) \
95+
hipLaunchKernelGGL(HIP_KERNEL_NAME(GPUCA_M_CAT(krnl_, GPUCA_M_KRNL_NAME(x_class))), dim3(x.nBlocks), dim3(x.nThreads), 0, me->mInternals->Streams[x.stream], GPUCA_CONSMEM_CALL y.index, args...);
96+
#endif // __HIPCC__
97+
98+
#include "GPUReconstructionKernelList.h"
99+
#undef GPUCA_KRNL
98100
#endif // ---------- COMPILE_MODE = onefile | rdc ----------
99-
100-
#include "GPUReconstructionKernelList.h"
101-
#undef GPUCA_KRNL
101+
// clang-format on
102102

103103
#ifndef GPUCA_NO_CONSTANT_MEMORY
104104
static GPUReconstructionDeviceBase::deviceConstantMemRegistration registerConstSymbol([]() {

‎GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAKernelsSpecialize.inc‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -97,7 +97,7 @@ struct GPUTPCGMO2OutputSort_comp {
9797
} // namespace o2::gpu::internal
9898

9999
template <>
100-
inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendInternal<GPUTPCGMMergerMergeBorders, 3>(const krnlSetupTime& _xyz, GPUTPCGMBorderRange* const& range, int32_t const& N, int32_t const& cmpMax)
100+
inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendTimed<GPUTPCGMMergerMergeBorders, 3>(const krnlSetupTime& _xyz, GPUTPCGMBorderRange* const& range, int32_t const& N, int32_t const& cmpMax)
101101
{
102102
if (cmpMax) {
103103
GPUCommonAlgorithm::sortOnDevice(this, _xyz.x.stream, range, N, MergeBorderTracks_compMax());
@@ -107,32 +107,32 @@ inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendInter
107107
}
108108

109109
template <>
110-
inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendInternal<GPUTPCGMMergerSortTracks, 0>(const krnlSetupTime& _xyz)
110+
inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendTimed<GPUTPCGMMergerSortTracks, 0>(const krnlSetupTime& _xyz)
111111
{
112112
GPUCommonAlgorithm::sortOnDevice(this, _xyz.x.stream, mProcessorsShadow->tpcMerger.TrackOrderProcess(), processors()->tpcMerger.NOutputTracks(), GPUTPCGMMergerSortTracks_comp(mProcessorsShadow->tpcMerger.OutputTracks()));
113113
}
114114

115115
template <>
116-
inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendInternal<GPUTPCGMMergerSortTracksQPt, 0>(const krnlSetupTime& _xyz)
116+
inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendTimed<GPUTPCGMMergerSortTracksQPt, 0>(const krnlSetupTime& _xyz)
117117
{
118118
GPUCommonAlgorithm::sortOnDevice(this, _xyz.x.stream, mProcessorsShadow->tpcMerger.TrackSort(), processors()->tpcMerger.NOutputTracks(), GPUTPCGMMergerSortTracksQPt_comp(mProcessorsShadow->tpcMerger.OutputTracks()));
119119
}
120120

121121
template <>
122-
inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendInternal<GPUTPCGMMergerMergeLoopers, 1>(const krnlSetupTime& _xyz)
122+
inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendTimed<GPUTPCGMMergerMergeLoopers, 1>(const krnlSetupTime& _xyz)
123123
{
124124
GPUCommonAlgorithm::sortOnDevice(this, _xyz.x.stream, mProcessorsShadow->tpcMerger.LooperCandidates(), processors()->tpcMerger.Memory()->nLooperMatchCandidates, GPUTPCGMMergerMergeLoopers_comp());
125125
}
126126

127127
template <>
128-
inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendInternal<GPUTPCGMO2Output, GPUTPCGMO2Output::sort>(const krnlSetupTime& _xyz)
128+
inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendTimed<GPUTPCGMO2Output, GPUTPCGMO2Output::sort>(const krnlSetupTime& _xyz)
129129
{
130130
GPUCommonAlgorithm::sortOnDevice(this, _xyz.x.stream, mProcessorsShadow->tpcMerger.TrackSortO2(), processors()->tpcMerger.NOutputTracksTPCO2(), GPUTPCGMO2OutputSort_comp());
131131
}
132132
#endif // GPUCA_SPECIALIZE_THRUST_SORTS
133133

134134
template <>
135-
inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendInternal<GPUMemClean16, 0>(const krnlSetupTime& _xyz, void* const& ptr, uint64_t const& size)
135+
inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendTimed<GPUMemClean16, 0>(const krnlSetupTime& _xyz, void* const& ptr, uint64_t const& size)
136136
{
137137
GPUChkErr(cudaMemsetAsync(ptr, 0, size, mInternals->Streams[_xyz.x.stream]));
138138
}

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -13,8 +13,8 @@
1313
/// \author David Rohr
1414

1515
#include "GPUReconstructionOCLIncludesHost.h"
16-
#include "GPUReconstructionProcessingKernels.inc"
1716
#include "GPUDefParametersLoad.inc"
17+
#include "GPUConstantMem.h"
1818

1919
#include <map>
2020

‎GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.h‎

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -34,7 +34,7 @@ class GPUReconstructionOCL : public GPUReconstructionProcessing::KernelInterface
3434
~GPUReconstructionOCL() override;
3535

3636
template <class T, int32_t I = 0, typename... Args>
37-
void runKernelBackend(const krnlSetupArgs<T, I, Args...>& args);
37+
void runKernelBackend(const krnlSetupTime& _xyz, const Args&... args);
3838

3939
protected:
4040
int32_t InitDevice_Runtime() override;
@@ -57,8 +57,6 @@ class GPUReconstructionOCL : public GPUReconstructionProcessing::KernelInterface
5757

5858
template <class T, int32_t I = 0>
5959
int32_t AddKernel();
60-
template <class T, int32_t I = 0, typename... Args>
61-
void runKernelBackendInternal(const krnlSetupTime& _xyz, const Args&... args);
6260

6361
GPUReconstructionOCLInternals* mInternals;
6462
float mOclVersion;

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

Lines changed: 3 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -16,9 +16,11 @@
1616
#include "GPUReconstructionKernelIncludes.h"
1717

1818
#include "GPUReconstructionOCLKernelsSpecialize.inc"
19+
#include "GPUReconstructionProcessingKernels.inc"
20+
template void GPUReconstructionProcessing::KernelInterface<GPUReconstructionOCL, GPUReconstructionDeviceBase>::runKernelVirtual(const int num, const void* args);
1921

2022
template <class T, int32_t I, typename... Args>
21-
inline void GPUReconstructionOCL::runKernelBackendInternal(const krnlSetupTime& _xyz, const Args&... args)
23+
inline void GPUReconstructionOCL::runKernelBackend(const krnlSetupTime& _xyz, const Args&... args)
2224
{
2325
cl_kernel k = getKernelObject<cl_kernel, T, I>();
2426
auto& x = _xyz.x;
@@ -48,12 +50,6 @@ inline void GPUReconstructionOCL::runKernelBackendInternal(const krnlSetupTime&
4850
}
4951
}
5052

51-
template <class T, int32_t I, typename... Args>
52-
void GPUReconstructionOCL::runKernelBackend(const krnlSetupArgs<T, I, Args...>& args)
53-
{
54-
std::apply([this, &args](auto&... vals) { runKernelBackendInternal<T, I, Args...>(args.s, vals...); }, args.v);
55-
}
56-
5753
template <class T, int32_t I>
5854
int32_t GPUReconstructionOCL::AddKernel()
5955
{
@@ -86,7 +82,3 @@ int32_t GPUReconstructionOCL::AddKernels()
8682
#undef GPUCA_KRNL
8783
return 0;
8884
}
89-
90-
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward, x_types, ...) template void GPUReconstructionOCL::runKernelBackend<GPUCA_M_KRNL_TEMPLATE(x_class)>(const krnlSetupArgs<GPUCA_M_KRNL_TEMPLATE(x_class) GPUCA_M_STRIP(x_types)>& args);
91-
#include "GPUReconstructionKernelList.h"
92-
#undef GPUCA_KRNL

0 commit comments

Comments
 (0)