GPU: Get rid of GPUCA_RTC_LB_..., use only GPUCA_LB_...

Author: davidrohr

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.h

@@ -45,8 +45,6 @@ class GPUReconstructionCUDABackend : public GPUReconstructionDeviceBase
   template <class T, int32_t I = 0, typename... Args>
   void runKernelBackendInternal(const krnlSetupTime& _xyz, const Args&... args);
 
-  void getRTCKernelCalls(std::vector<std::string>& kernels);
-
   template <class T, class S>
   friend GPUh() void GPUCommonAlgorithm::sortOnDevice(auto* rec, int32_t stream, T* begin, size_t N, const S& comp);
   GPUReconstructionCUDAInternals* mInternals;
@@ -91,6 +89,7 @@ class GPUReconstructionCUDA : public GPUReconstructionKernels<GPUReconstructionC
 
  private:
   int32_t genRTC(std::string& filename, uint32_t& nCompile);
+  void getRTCKernelCalls(std::vector<std::string>& kernels);
   void genAndLoadRTC();
   void loadKernelModules(bool perKernel);
   const char *mRtcSrcExtension = ".src", *mRtcBinExtension = ".o";

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAGenRTC.cxx

@@ -272,3 +272,13 @@ int32_t GPUReconstructionCUDA::genRTC(std::string& filename, uint32_t& nCompile)
 
   return 0;
 }
+
+void GPUReconstructionCUDA::getRTCKernelCalls(std::vector<std::string>& kernels)
+{
+#undef GPUCA_KRNL_LB
+#define GPUCA_KRNL(...) kernels.emplace_back(GPUCA_M_STR(GPUCA_KRNLGPU(__VA_ARGS__)));
+#define GPUCA_KRNL_LB(x_class, x_attributes, ...) GPUCA_KRNL(x_class, (REG, (GPUCA_M_CAT(GPUCA_LB_, GPUCA_M_KRNL_NAME(x_class))), GPUCA_M_STRIP(x_attributes)), __VA_ARGS__)
+#include "GPUReconstructionKernelList.h"
+#undef GPUCA_KRNL
+#undef GPUCA_KRNL_LB
+}

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAIncludesSystem.h

@@ -12,8 +12,8 @@
 /// \file GPUReconstructionCUDAIncludesSystem.h
 /// \author David Rohr
 
-#ifndef O2_GPU_GPURECONSTRUCTIONCUDAINCLUDES_H
-#define O2_GPU_GPURECONSTRUCTIONCUDAINCLUDES_H
+#ifndef O2_GPU_GPURECONSTRUCTIONCUDAINCLUDESSYSTEM_H
+#define O2_GPU_GPURECONSTRUCTIONCUDAINCLUDESSYSTEM_H
 
 #include <cstdint>
 #include <type_traits>
@@ -32,4 +32,4 @@
 #include <sm_20_atomic_functions.h>
 #include <cuda_fp16.h>
 
-#endif
+#endif // O2_GPU_GPURECONSTRUCTIONCUDAINCLUDESSYSTEM_H

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAKernels.cu

@@ -118,14 +118,3 @@ static GPUReconstructionDeviceBase::deviceConstantMemRegistration registerConstS
   return retVal;
 });
 #endif
-
-void GPUReconstructionCUDABackend::getRTCKernelCalls(std::vector<std::string>& kernels)
-{
-#undef GPUCA_KRNL_LB
-#undef __launch_bounds__
-#define GPUCA_KRNL(...) kernels.emplace_back(GPUCA_M_STR(GPUCA_KRNLGPU(__VA_ARGS__)));
-#define GPUCA_KRNL_LB(x_class, x_attributes, ...) GPUCA_KRNL(x_class, (REG, (GPUCA_M_CAT(GPUCA_RTC_LB_, GPUCA_M_KRNL_NAME(x_class))), GPUCA_M_STRIP(x_attributes)), __VA_ARGS__)
-#include "GPUReconstructionKernelList.h"
-#undef GPUCA_KRNL
-#undef GPUCA_KRNL_LB
-}

GPU/GPUTracking/Base/hip/GPUReconstructionHIPIncludesSystem.h

@@ -12,8 +12,8 @@
 /// \file GPUReconstructionHIPIncludesSystem.h
 /// \author David Rohr
 
-#ifndef O2_GPU_RECONSTRUCTIONHIPINCLUDES_H
-#define O2_GPU_RECONSTRUCTIONHIPINCLUDES_H
+#ifndef O2_GPU_RECONSTRUCTIONHIPINCLUDESSYSTEM_H
+#define O2_GPU_RECONSTRUCTIONHIPINCLUDESSYSTEM_H
 
 #include <hip/hip_runtime.h>
 #include <hip/hip_ext.h>
@@ -25,4 +25,4 @@
 #include <thrust/device_ptr.h>
 #pragma GCC diagnostic pop
 
-#endif
+#endif // O2_GPU_RECONSTRUCTIONHIPINCLUDESSYSTEM_H

GPU/GPUTracking/Definitions/GPUDefParametersLoad.template.inc

@@ -37,17 +37,17 @@ static GPUDefParameters GPUDefParametersLoad()
   };
 }
 
-#define GPUCA_EXPORT_KERNEL(name)                                                                                  \
-  if (par.par_LB_maxThreads[i] > 0) {                                                                              \
-    o << "#define GPUCA_" << (forRTC ? "RTC_" : "") << "LB_" << GPUCA_M_STR(name) " " << par.par_LB_maxThreads[i]; \
-    if (par.par_LB_minBlocks[i] > 0) {                                                                             \
-      o << ", " << par.par_LB_minBlocks[i];                                                                        \
-    }                                                                                                              \
-    if (!forRTC && par.par_LB_forceBlocks[i] > 0) {                                                                \
-      o << ", " << par.par_LB_forceBlocks[i];                                                                      \
-    }                                                                                                              \
-    o << "\n";                                                                                                     \
-  }                                                                                                                \
+#define GPUCA_EXPORT_KERNEL(name)                                                  \
+  if (par.par_LB_maxThreads[i] > 0) {                                              \
+    o << "#define GPUCA_LB_" << GPUCA_M_STR(name) " " << par.par_LB_maxThreads[i]; \
+    if (par.par_LB_minBlocks[i] > 0) {                                             \
+      o << ", " << par.par_LB_minBlocks[i];                                        \
+    }                                                                              \
+    if (!forRTC && par.par_LB_forceBlocks[i] > 0) {                                \
+      o << ", " << par.par_LB_forceBlocks[i];                                      \
+    }                                                                              \
+    o << "\n";                                                                     \
+  }                                                                                \
   i++;
 
 static std::string GPUDefParametersExport(const GPUDefParameters& par, bool forRTC)

GPU/GPUTracking/Definitions/GPUDefParametersWrapper.h

@@ -22,7 +22,7 @@
 #include "GPUCommonDef.h"
 #include "GPUDefMacros.h"
 
-#ifndef GPUCA_GPUCODE_GENRTC
+#if defined(GPUCA_GPUCODE) && !defined(GPUCA_GPUCODE_GENRTC)
 #include "GPUDefParametersDefaults.h"
 #endif
 #include "GPUDefParametersConstants.h"