GPU: Cleanup unused template parameter

davidrohr · davidrohr · commit f26e72529547 · 2025-03-26T19:59:45.000+01:00
diff --git a/GPU/GPUTracking/Base/GPUReconstruction.cxx b/GPU/GPUTracking/Base/GPUReconstruction.cxx
@@ -273,6 +273,7 @@ int32_t GPUReconstruction::InitPhaseBeforeDevice()
     if (mProcessingSettings.createO2Output > 1) {
       mProcessingSettings.createO2Output = 1;
     }
+    mProcessingSettings.rtc.deterministic = 1;
   }
   if (mProcessingSettings.deterministicGPUReconstruction && mProcessingSettings.debugLevel >= 6) {
     mProcessingSettings.nTPCClustererLanes = 1;
diff --git a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu
@@ -618,7 +618,7 @@ void GPUReconstructionCUDA::loadKernelModules(bool perKernel)
 {
   uint32_t j = 0;
 #define GPUCA_KRNL(x_class, ...)                                                                                                                                                        \
-  getRTCkernelNum<false, GPUCA_M_KRNL_TEMPLATE(x_class)>(mInternals->kernelFunctions.size());                                                                                           \
+  getRTCkernelNum<GPUCA_M_KRNL_TEMPLATE(x_class)>(mInternals->kernelFunctions.size());                                                                                                  \
   mInternals->kernelFunctions.emplace_back(new CUfunction);                                                                                                                             \
   mInternals->kernelNames.emplace_back(GPUCA_M_STR(GPUCA_M_CAT(krnl_, GPUCA_M_KRNL_NAME(x_class))));                                                                                    \
   if (mProcessingSettings.debugLevel >= 3) {                                                                                                                                            \
diff --git a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.h b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.h
@@ -49,7 +49,7 @@ class GPUReconstructionCUDABackend : public GPUReconstructionDeviceBase
   template <class T, int32_t I>
   class backendInternal;
 
-  template <bool multi, class T, int32_t I = 0>
+  template <class T, int32_t I = 0>
   static int32_t getRTCkernelNum(int32_t k = -1);
 
   void getRTCKernelCalls(std::vector<std::string>& kernels);
diff --git a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAKernels.cu b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAKernels.cu
@@ -55,7 +55,7 @@ inline void GPUReconstructionCUDABackend::runKernelBackendInternal(const krnlSet
 #endif
     pArgs[arg_offset] = &y.index;
     GPUReconstructionCUDAInternals::getArgPtrs(&pArgs[arg_offset + 1], args...);
-    GPUChkErr(cuLaunchKernel(*mInternals->kernelFunctions[getRTCkernelNum<false, T, I>()], x.nBlocks, 1, 1, x.nThreads, 1, 1, 0, mInternals->Streams[x.stream], (void**)pArgs, nullptr));
+    GPUChkErr(cuLaunchKernel(*mInternals->kernelFunctions[getRTCkernelNum<T, I>()], x.nBlocks, 1, 1, x.nThreads, 1, 1, 0, mInternals->Streams[x.stream], (void**)pArgs, nullptr));
   }
 }
 
@@ -111,7 +111,7 @@ void GPUReconstructionCUDABackend::runKernelBackend(const krnlSetupArgs<T, I, Ar
 #include "GPUReconstructionKernelList.h"
 #undef GPUCA_KRNL
 
-template <bool multi, class T, int32_t I>
+template <class T, int32_t I>
 int32_t GPUReconstructionCUDABackend::getRTCkernelNum(int32_t k)
 {
   static int32_t num = k;
@@ -121,9 +121,7 @@ int32_t GPUReconstructionCUDABackend::getRTCkernelNum(int32_t k)
   return num;
 }
 
-#define GPUCA_KRNL(x_class, ...)                                                                                    \
-  template int32_t GPUReconstructionCUDABackend::getRTCkernelNum<false, GPUCA_M_KRNL_TEMPLATE(x_class)>(int32_t k); \
-  template int32_t GPUReconstructionCUDABackend::getRTCkernelNum<true, GPUCA_M_KRNL_TEMPLATE(x_class)>(int32_t k);
+#define GPUCA_KRNL(x_class, ...) template int32_t GPUReconstructionCUDABackend::getRTCkernelNum<GPUCA_M_KRNL_TEMPLATE(x_class)>(int32_t k);
 #include "GPUReconstructionKernelList.h"
 #undef GPUCA_KRNL
 

Original file line number	Diff line number	Diff line change
`@@ -273,6 +273,7 @@ int32_t GPUReconstruction::InitPhaseBeforeDevice()`
`273`	`273`	`if (mProcessingSettings.createO2Output > 1) {`
`274`	`274`	`mProcessingSettings.createO2Output = 1;`
`275`	`275`	`}`
	`276`	`+ mProcessingSettings.rtc.deterministic = 1;`
`276`	`277`	`}`
`277`	`278`	`if (mProcessingSettings.deterministicGPUReconstruction && mProcessingSettings.debugLevel >= 6) {`
`278`	`279`	`mProcessingSettings.nTPCClustererLanes = 1;`
Original file line number	Diff line number	Diff line change
`@@ -618,7 +618,7 @@ void GPUReconstructionCUDA::loadKernelModules(bool perKernel)`
`618`	`618`	`{`
`619`	`619`	`uint32_t j = 0;`
`620`	`620`	`#define GPUCA_KRNL(x_class, ...) \`
`621`		`- getRTCkernelNum<false, GPUCA_M_KRNL_TEMPLATE(x_class)>(mInternals->kernelFunctions.size()); \`
	`621`	`+ getRTCkernelNum<GPUCA_M_KRNL_TEMPLATE(x_class)>(mInternals->kernelFunctions.size()); \`
`622`	`622`	`mInternals->kernelFunctions.emplace_back(new CUfunction); \`
`623`	`623`	`mInternals->kernelNames.emplace_back(GPUCA_M_STR(GPUCA_M_CAT(krnl_, GPUCA_M_KRNL_NAME(x_class)))); \`
`624`	`624`	`if (mProcessingSettings.debugLevel >= 3) { \`
Original file line number	Diff line number	Diff line change
`@@ -55,7 +55,7 @@ inline void GPUReconstructionCUDABackend::runKernelBackendInternal(const krnlSet`
`55`	`55`	`#endif`
`56`	`56`	`pArgs[arg_offset] = &y.index;`
`57`	`57`	`GPUReconstructionCUDAInternals::getArgPtrs(&pArgs[arg_offset + 1], args...);`
`58`		`- GPUChkErr(cuLaunchKernel(mInternals->kernelFunctions[getRTCkernelNum<false, T, I>()], x.nBlocks, 1, 1, x.nThreads, 1, 1, 0, mInternals->Streams[x.stream], (void*)pArgs, nullptr));`
	`58`	`+ GPUChkErr(cuLaunchKernel(mInternals->kernelFunctions[getRTCkernelNum<T, I>()], x.nBlocks, 1, 1, x.nThreads, 1, 1, 0, mInternals->Streams[x.stream], (void*)pArgs, nullptr));`
`59`	`59`	`}`
`60`	`60`	`}`
`61`	`61`
`@@ -111,7 +111,7 @@ void GPUReconstructionCUDABackend::runKernelBackend(const krnlSetupArgs<T, I, Ar`
`111`	`111`	`#include "GPUReconstructionKernelList.h"`
`112`	`112`	`#undef GPUCA_KRNL`
`113`	`113`
`114`		`-template <bool multi, class T, int32_t I>`
	`114`	`+template <class T, int32_t I>`
`115`	`115`	`int32_t GPUReconstructionCUDABackend::getRTCkernelNum(int32_t k)`
`116`	`116`	`{`
`117`	`117`	`static int32_t num = k;`
`@@ -121,9 +121,7 @@ int32_t GPUReconstructionCUDABackend::getRTCkernelNum(int32_t k)`
`121`	`121`	`return num;`
`122`	`122`	`}`
`123`	`123`
`124`		`-#define GPUCA_KRNL(x_class, ...) \`
`125`		`- template int32_t GPUReconstructionCUDABackend::getRTCkernelNum<false, GPUCA_M_KRNL_TEMPLATE(x_class)>(int32_t k); \`
`126`		`- template int32_t GPUReconstructionCUDABackend::getRTCkernelNum<true, GPUCA_M_KRNL_TEMPLATE(x_class)>(int32_t k);`
	`124`	`+#define GPUCA_KRNL(x_class, ...) template int32_t GPUReconstructionCUDABackend::getRTCkernelNum<GPUCA_M_KRNL_TEMPLATE(x_class)>(int32_t k);`
`127`	`125`	`#include "GPUReconstructionKernelList.h"`
`128`	`126`	`#undef GPUCA_KRNL`
`129`	`127`