GPU: Make GPUDefParamters available as C++ struct at runtime

Author: davidrohr

GPU/GPUTracking/Base/GPUReconstructionProcessing.cxx

@@ -14,9 +14,31 @@
 
 #include "GPUReconstructionProcessing.h"
 #include "GPUReconstructionThreading.h"
+#include "GPUDefParametersDefault.h"
+#include "GPUDefParametersLoad.inc"
 
 using namespace o2::gpu;
 
+GPUReconstructionProcessing::GPUReconstructionProcessing(const GPUSettingsDeviceBackend& cfg) : GPUReconstruction(cfg)
+{
+  if (mMaster == nullptr) {
+    mParCPU = new GPUDefParameters(o2::gpu::internal::GPUDefParametersLoad());
+    mParDevice = new GPUDefParameters();
+  } else {
+    GPUReconstructionProcessing* master = dynamic_cast<GPUReconstructionProcessing*>(mMaster);
+    mParCPU = master->mParCPU;
+    mParDevice = master->mParDevice;
+  }
+}
+
+GPUReconstructionProcessing::~GPUReconstructionProcessing()
+{
+  if (mMaster == nullptr) {
+    delete mParCPU;
+    delete mParDevice;
+  }
+}
+
 int32_t GPUReconstructionProcessing::getNKernelHostThreads(bool splitCores)
 {
   int32_t nThreads = 0;

GPU/GPUTracking/Base/GPUReconstructionProcessing.h

@@ -25,6 +25,8 @@
 namespace o2::gpu
 {
 
+struct GPUDefParameters;
+
 namespace gpu_reconstruction_kernels
 {
 struct deviceEvent {
@@ -63,7 +65,7 @@ class threadContext
 class GPUReconstructionProcessing : public GPUReconstruction
 {
  public:
-  ~GPUReconstructionProcessing() override = default;
+  ~GPUReconstructionProcessing() override;
 
   // Threading
   int32_t getNKernelHostThreads(bool splitCores);
@@ -101,7 +103,7 @@ class GPUReconstructionProcessing : public GPUReconstruction
   };
 
  protected:
-  GPUReconstructionProcessing(const GPUSettingsDeviceBackend& cfg) : GPUReconstruction(cfg) {}
+  GPUReconstructionProcessing(const GPUSettingsDeviceBackend& cfg);
   using deviceEvent = gpu_reconstruction_kernels::deviceEvent;
 
   static const std::vector<std::string> mKernelNames;
@@ -132,6 +134,9 @@ class GPUReconstructionProcessing : public GPUReconstruction
   template <class T, int32_t J = -1>
   HighResTimer& getTimer(const char* name, int32_t num = -1);
 
+  GPUDefParameters* mParCPU = nullptr;
+  GPUDefParameters* mParDevice = nullptr;
+
  private:
   uint32_t getNextTimerId();
   timerMeta* getTimerById(uint32_t id, bool increment = true);

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu

@@ -47,15 +47,21 @@ __global__ void dummyInitKernel(void*) {}
 
 #include "GPUReconstructionIncludesITS.h"
 
+#include "GPUDefParametersDefault.h"
+#include "GPUDefParametersLoad.inc"
+
 GPUReconstructionCUDABackend::GPUReconstructionCUDABackend(const GPUSettingsDeviceBackend& cfg) : GPUReconstructionDeviceBase(cfg, sizeof(GPUReconstructionDeviceBase))
 {
   if (mMaster == nullptr) {
     mInternals = new GPUReconstructionCUDAInternals;
+    *mParDevice = o2::gpu::internal::GPUDefParametersLoad();
   }
+  mDeviceBackendSettings.deviceType = DeviceType::CUDA;
 }
 
 GPUReconstructionCUDABackend::~GPUReconstructionCUDABackend()
 {
+  Exit(); // Make sure we destroy everything (in particular the ITS tracker) before we exit CUDA
   if (mMaster == nullptr) {
     delete mInternals;
   }
@@ -69,7 +75,6 @@ int32_t GPUReconstructionCUDABackend::GPUChkErrInternal(const int64_t error, con
 
 GPUReconstructionCUDA::GPUReconstructionCUDA(const GPUSettingsDeviceBackend& cfg) : GPUReconstructionKernels(cfg)
 {
-  mDeviceBackendSettings.deviceType = DeviceType::CUDA;
 #ifndef __HIPCC__ // CUDA
   mRtcSrcExtension = ".cu";
   mRtcBinExtension = ".fatbin";
@@ -78,11 +83,7 @@ GPUReconstructionCUDA::GPUReconstructionCUDA(const GPUSettingsDeviceBackend& cfg
   mRtcBinExtension = ".o";
 #endif
 }
-
-GPUReconstructionCUDA::~GPUReconstructionCUDA()
-{
-  Exit(); // Make sure we destroy everything (in particular the ITS tracker) before we exit CUDA
-}
+GPUReconstructionCUDA::~GPUReconstructionCUDA() {}
 
 GPUReconstruction* GPUReconstruction_Create_CUDA(const GPUSettingsDeviceBackend& cfg) { return new GPUReconstructionCUDA(cfg); }
 

GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.cxx

@@ -14,6 +14,9 @@
 
 #include "GPUReconstructionOCLIncludesHost.h"
 
+#include "GPUDefParametersDefault.h"
+#include "GPUDefParametersLoad.inc"
+
 #include <map>
 
 static_assert(std::is_convertible<cl_event, void*>::value, "OpenCL event type incompatible to deviceEvent");
@@ -36,6 +39,7 @@ GPUReconstructionOCLBackend::GPUReconstructionOCLBackend(const GPUSettingsDevice
 {
   if (mMaster == nullptr) {
     mInternals = new GPUReconstructionOCLInternals;
+    *mParDevice = o2::gpu::internal::GPUDefParametersLoad();
   }
   mDeviceBackendSettings.deviceType = DeviceType::OCL;
 }

GPU/GPUTracking/CMakeLists.txt

@@ -128,7 +128,7 @@ set(HDRS_INSTALL
     DataTypes/GPUTriggerOutputs.h
     Debug/GPUROOTDump.h
     Definitions/GPUDefConstantsAndSettings.h
-    Definitions/GPUDefGPUParameters.h
+    Definitions/GPUDefParametersDefault.h
     Definitions/GPUDef.h
     Definitions/GPUDefMacros.h
     Definitions/GPULogging.h
@@ -246,10 +246,18 @@ file(GENERATE
      OUTPUT include_gpu_onthefly/GPUNoFastMathKernels.h
      INPUT cmake/GPUNoFastMathKernels.template.h
 )
+file(GENERATE
+     OUTPUT include_gpu_onthefly/GPUDefParameters.h
+     INPUT Definitions/GPUDefParameters.template.h
+)
+file(GENERATE
+     OUTPUT include_gpu_onthefly/GPUDefParametersLoad.inc
+     INPUT Definitions/GPUDefParametersLoad.template.inc
+)
 if(NOT ALIGPU_BUILD_TYPE STREQUAL "O2")
   include_directories(${CMAKE_CURRENT_BINARY_DIR}/include_gpu_onthefly)
 endif()
-set(HDRS_INSTALL ${HDRS_INSTALL} ${CMAKE_CURRENT_BINARY_DIR}/include_gpu_onthefly/GPUReconstructionKernelList.h)
+set(HDRS_INSTALL ${HDRS_INSTALL} ${CMAKE_CURRENT_BINARY_DIR}/include_gpu_onthefly/GPUReconstructionKernelList.h ${CMAKE_CURRENT_BINARY_DIR}/include_gpu_onthefly/GPUDefParameters.h)
 include(kernels.cmake)
 
 # Optional sources depending on optional dependencies

GPU/GPUTracking/Definitions/GPUDef.h

@@ -18,7 +18,7 @@
 
 #include "GPUCommonDef.h"
 #include "GPUDefConstantsAndSettings.h"
-#include "GPUDefGPUParameters.h"
+#include "GPUDefParametersDefault.h"
 #include "GPUCommonRtypes.h"
 
 // Macros for masking ptrs in OpenCL kernel calls as uint64_t (The API only allows us to pass buffer objects)

GPU/GPUTracking/Definitions/GPUDefConstantsAndSettings.h

@@ -13,7 +13,7 @@
 /// \author David Rohr
 
 // This files contains compile-time constants affecting the GPU algorithms / reconstruction results.
-// Architecture-dependant compile-time constants affecting the performance without changing the results are stored in GPUDefGPUParameters.h
+// Architecture-dependant compile-time constants affecting the performance without changing the results are stored in GPUDefParameters.h
 
 #ifndef GPUDEFCONSTANTSANDSETTINGS_H
 #define GPUDEFCONSTANTSANDSETTINGS_H

GPU/GPUTracking/Definitions/GPUDefParameters.template.h

@@ -0,0 +1,28 @@
+// Copyright 2019-2020 CERN and copyright holders of ALICE O2.
+// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
+// All rights not expressly granted are reserved.
+//
+// This software is distributed under the terms of the GNU General Public
+// License v3 (GPL Version 3), copied verbatim in the file "COPYING".
+//
+// In applying this license CERN does not waive the privileges and immunities
+// granted to it by virtue of its status as an Intergovernmental Organization
+// or submit itself to any jurisdiction.
+
+/// \file GPUDefParameters.h
+/// \author David Rohr
+
+#ifndef GPUDEFPARAMETERS_H
+#define GPUDEFPARAMETERS_H
+
+namespace o2::gpu
+{
+struct GPUDefParameters {  // clang-format off
+  $<JOIN:$<LIST:TRANSFORM,$<LIST:TRANSFORM,$<LIST:TRANSFORM,$<TARGET_PROPERTY:O2_GPU_KERNELS,O2_GPU_KERNEL_NAMES>,REPLACE,[^A-Za-z0-9]+,_>,PREPEND,int32_t parLB_>,APPEND,_1 = -1>,;
+  >;
+  $<JOIN:$<LIST:TRANSFORM,$<LIST:TRANSFORM,$<LIST:TRANSFORM,$<TARGET_PROPERTY:O2_GPU_KERNELS,O2_GPU_KERNEL_NAMES>,REPLACE,[^A-Za-z0-9]+,_>,PREPEND,int32_t parLB_>,APPEND,_2 = -1>,;
+  >;
+};  // clang-format on
+}  // namespace o2::gpu
+
+#endif