GPUDisplay: Many improvements, implement Qt GUI and Configuration Console, many fixes and cleanups

.clang-format

@@ -27,6 +27,7 @@ IndentWidth:     2
 # It is broken on windows. Breaks all #include "header.h"
 ---
 Language:        Cpp
+AttributeMacros: ['GPUrestrict'] # GPU qualifier macro: GPUrestrict() is otherwise parsed as function call, e.g. 'a -> b' after it
 MaxEmptyLinesToKeep: 1
 KeepEmptyLinesAtTheStartOfBlocks: true
 NamespaceIndentation: None

GPU/Common/GPUCommonArray.h

@@ -41,8 +41,10 @@ struct array {
   }
   T m_internal_V__[N];
 };
+// clang-format off
 template <class T, class... E>
 GPUd() array(T, E...)->array<T, 1 + sizeof...(E)>;
+// clang-format on
 #else
 template <typename T, size_t N>
 using array = std::array<T, N>;

GPU/Common/GPUCommonConstants.h

@@ -21,6 +21,6 @@ namespace o2::gpu::gpu_common_constants
 {
 static GPUglobalconstexpr() const float kCLight = 0.000299792458f; // TODO: Duplicate of MathConstants, fix this now that we use only OpenCL CPP
 static GPUglobalconstexpr() const float kZeroFieldCut = 0.013f;
-}
+} // namespace o2::gpu::gpu_common_constants
 
 #endif

GPU/GPUTracking/Base/GPUParam.h

@@ -18,6 +18,7 @@
 #include "GPUCommonDef.h"
 #include "GPUCommonMath.h"
 #include "GPUDef.h"
+#define GPUSETTINGS_PARAM
 #include "GPUSettings.h"
 #include "GPUTPCGMPolynomialField.h"
 #include "DataFormatsTPC/Constants.h"

GPU/GPUTracking/Base/GPUParamRTC.h

@@ -16,6 +16,8 @@
 #define GPUPARAMRTC_H
 
 #include "GPUParam.h"
+#define GPUSETTINGS_ALL
+#include "GPUSettings.h"
 #include <string>
 
 namespace o2::gpu

GPU/GPUTracking/Base/GPUReconstructionKernelIncludes.template.h

@@ -14,5 +14,4 @@
 
 // clang-format off
 $<JOIN:$<LIST:TRANSFORM,$<LIST:TRANSFORM,$<LIST:REMOVE_DUPLICATES,$<TARGET_PROPERTY:O2_GPU_KERNELS,O2_GPU_KERNEL_INCLUDES>>,APPEND,.h">,PREPEND,#include ">,
->
-// clang-format on
+> // clang-format on

GPU/GPUTracking/Base/GPUReconstructionKernelList.template.h

@@ -16,5 +16,4 @@
 #include "GPUReconstructionKernelMacros.h"
 
 // clang-format off
-$<JOIN:$<TARGET_PROPERTY:O2_GPU_KERNELS,O2_GPU_KERNELS>,>
-// clang-format on
+$<JOIN:$<TARGET_PROPERTY:O2_GPU_KERNELS,O2_GPU_KERNELS>,> // clang-format on

GPU/GPUTracking/Base/GPUReconstructionTimeframe.cxx

@@ -24,6 +24,8 @@
 #include "AliHLTTPCRawCluster.h"
 #include "TPCFastTransformPOD.h"
 #include "GPUO2DataTypes.h"
+#define GPUSETTINGS_STANDALONE
+#define GPUSETTINGS_DISPLAY
 #include "GPUSettings.h"
 
 #include <cstdio>

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu

@@ -331,9 +331,9 @@ int32_t GPUReconstructionCUDA::InitDevice_Runtime()
       }
     }
 
-#ifndef __HIPCC__ // CUDA
+#ifndef __HIPCC__                                                      // CUDA
     dummyInitKernel<<<mMultiprocessorCount, 256>>>(mDeviceMemoryBase); // TODO: Can't we just use the CUDA version and hipify will take care of the rest?
-#else // HIP
+#else                                                                  // HIP
     hipLaunchKernelGGL(HIP_KERNEL_NAME(dummyInitKernel), dim3(mMultiprocessorCount), dim3(256), 0, 0, mDeviceMemoryBase);
 #endif
 

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAkernel.template.cu

@@ -21,10 +21,7 @@
 
 // clang-format off
 @O2_GPU_KERNEL_TEMPLATE_FILES@
-// clang-format on
 
 extern "C" {
-// clang-format off
 @O2_GPU_KERNEL_TEMPLATE_REPLACE@
-// clang-format on
-}
+} // clang-format on: not before the }, the placeholders confuse clang-format

GPU/GPUTracking/Base/hip/GPUReconstructionHIPkernel.template.hip

@@ -21,10 +21,7 @@
 
 // clang-format off
 @O2_GPU_KERNEL_TEMPLATE_FILES@
-// clang-format on
 
 extern "C" {
-// clang-format off
 @O2_GPU_KERNEL_TEMPLATE_REPLACE@
-// clang-format on
-}
+} // clang-format on: not before the }, the placeholders confuse clang-format

GPU/GPUTracking/CMakeLists.txt

@@ -69,7 +69,7 @@ set(SRCS_DATATYPES DataTypes/GPUDataTypesConfig.cxx DataTypes/GPUConfigDump.cxx
 
 set(HDRS_CINT_O2 Merger/GPUTPCGMTrackParam.h Merger/GPUTPCGMMergedTrack.h Merger/GPUTPCGMSectorTrack.h Merger/GPUTPCGMBorderTrack.h TRDTracking/GPUTRDInterfaces.h)
 set(HDRS_CINT_DATATYPES DataTypes/GPUTPCGMMergedTrackHit.h)
-set(HDRS_CINT_O2_ADDITIONAL DataTypes/GPUSettings.h Definitions/GPUSettingsList.h DataTypes/GPUDataTypesIO.h DataTypes/GPUDataTypesConfig.h DataTypes/GPUDataTypesQA.h DataTypes/GPUTRDTrack.h DataTypes/GPUTRDRecoParam.h DataTypes/CalibdEdxTrackTopologyPol.h DataTypes/CalibdEdxTrackTopologySpline.h) # Manual dependencies for ROOT dictionary generation
+set(HDRS_CINT_O2_ADDITIONAL DataTypes/GPUSettings.h DataTypes/GPUSettingsTypes.h Definitions/GPUSettingsListDisplay.h Definitions/GPUSettingsList.h DataTypes/GPUDataTypesIO.h DataTypes/GPUDataTypesConfig.h DataTypes/GPUDataTypesQA.h DataTypes/GPUTRDTrack.h DataTypes/GPUTRDRecoParam.h DataTypes/CalibdEdxTrackTopologyPol.h DataTypes/CalibdEdxTrackTopologySpline.h) # Manual dependencies for ROOT dictionary generation
 
 set(SRCS_NO_CINT
     DataTypes/GPUMemorySizeScalers.cxx
@@ -120,6 +120,7 @@ set(HDRS_INSTALL
     Base/GPUReconstructionIncludesITS.h
     Base/GPUReconstructionKernelMacros.h
     Base/GPUReconstructionCPUKernels.h
+    Base/GPUStdSystemHeaders.h
     DataCompression/GPUTPCClusterRejection.h
     DataCompression/GPUTPCCompressionKernels.inc
     DataCompression/TPCClusterDecompressionCore.inc
@@ -156,6 +157,7 @@ set(HDRS_INSTALL
     SectorTracker/GPUTPCMCInfo.h
     SectorTracker/GPUTPCTracklet.h
     SectorTracker/GPUTPCTrackLinearisation.h
+    SectorTracker/GPUTPCRow.h
     TPCConvert/GPUTPCConvertImpl.h
     TRDTracking/GPUTRDGeometry.h
     TRDTracking/GPUTRDSpacePoint.h
@@ -165,6 +167,8 @@ set(HDRS_INSTALL
     TRDTracking/GPUTRDTrackPoint.h
     TRDTracking/GPUTRDTrackPoint.h
     DataTypes/GPUTPCGMPolynomialField.h
+    DataTypes/GPUSettingsParam.h
+    Definitions/GPUSettingsListStandalone.h
 )
 
 set(SRCS_NO_CINT ${SRCS_NO_CINT} display/GPUDisplayInterface.cxx)

GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx

@@ -67,7 +67,7 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step0at
       if (!(param.rec.tpc.compressionTypeMask & GPUSettings::CompressionTrackModel)) {
         continue; // No track model compression
       }
-      const ClusterNative& GPUrestrict() orgCl = clusters -> clusters[hit.sector][hit.row][hit.num - clusters->clusterOffset[hit.sector][hit.row]];
+      const ClusterNative& GPUrestrict() orgCl = clusters->clusters[hit.sector][hit.row][hit.num - clusters->clusterOffset[hit.sector][hit.row]];
       constexpr GPUTPCGeometry geo;
       float x = geo.Row2X(hit.row);
       float y = track.LinearPad2Y(hit.sector, orgCl.getPad(), geo.PadWidth(hit.row), geo.NPads(hit.row));
@@ -291,7 +291,7 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step1un
 
       for (uint32_t j = get_local_id(0); j < count; j += get_local_size(0)) {
         int32_t outidx = idOffsetOut + totalCount + j;
-        const ClusterNative& GPUrestrict() orgCl = clusters -> clusters[iSector][iRow][sortBuffer[j]];
+        const ClusterNative& GPUrestrict() orgCl = clusters->clusters[iSector][iRow][sortBuffer[j]];
 
         int32_t preId = j != 0 ? (int32_t)sortBuffer[j - 1] : (totalCount != 0 ? (int32_t)smem.lastIndex : -1);
         GPUTPCCompression_EncodeUnattached(param.rec.tpc.compressionTypeMask, orgCl, c.timeDiffU[outidx], c.padDiffU[outidx], preId == -1 ? nullptr : &clusters->clusters[iSector][iRow][preId]);

GPU/GPUTracking/DataTypes/GPUConfigDump.cxx

@@ -14,6 +14,7 @@
 
 #include "GPUConfigDump.h"
 #include "GPUDataTypesIO.h"
+#define GPUSETTINGS_ALL
 #include "GPUSettings.h"
 
 #include <functional>
@@ -27,10 +28,7 @@ using namespace o2::gpu;
 namespace
 {
 GPUSettingsStandalone configStandalone;
-std::vector<std::function<void()>> qprint_global;
-#define QCONFIG_PRINT
-#include "utils/qconfig_internal.h"
-#undef QCONFIG_PRINT
+#include "utils/qconfig_print.inc"
 } // namespace
 
 void GPUConfigDump::dumpConfig(const GPUSettingsRec* rec, const GPUSettingsProcessing* proc, const GPUSettingsQA* qa, const GPUSettingsDisplay* display, const GPUSettingsDeviceBackend* device, const GPURecoStepConfiguration* workflow)

GPU/GPUTracking/DataTypes/GPUO2ConfigurableParam.cxx

@@ -35,6 +35,7 @@ using namespace o2::gpu;
 #define AddCustomCPP(...)
 #define AddHelp(...)
 #define AddShortcut(...)
+#include "GPUSettingsListDisplay.h"
 #include "GPUSettingsList.h"
 #undef BeginNamespace
 #undef EndNamespace
@@ -80,6 +81,7 @@ GPUSettingsO2 GPUO2InterfaceConfiguration::ReadConfigurableParam(GPUO2InterfaceC
 #define AddCustomCPP(...)
 #define AddHelp(...)
 #define AddShortcut(...)
+#include "GPUSettingsListDisplay.h"
 #include "GPUSettingsList.h"
 #undef BeginNamespace
 #undef EndNamespace

GPU/GPUTracking/DataTypes/GPUO2ConfigurableParam.h

@@ -25,6 +25,7 @@
 #include "GPUO2ExternalUser.h"
 #include "CommonUtils/ConfigurableParam.h"
 #include "CommonUtils/ConfigurableParamHelper.h"
+#define GPUSETTINGS_DISPLAY
 #include "GPUSettings.h"
 #include "GPUDefMacros.h"
 #include <vector>
@@ -59,6 +60,7 @@
 #define AddShortcut(...)
 #define AddOptionRTC(...) AddOption(__VA_ARGS__)
 #define AddOptionArrayRTC(...) AddOptionArray(__VA_ARGS__)
+#include "GPUSettingsListDisplay.h"
 #include "GPUSettingsList.h"
 #undef BeginNamespace
 #undef EndNamespace

GPU/GPUTracking/DataTypes/GPUSettings.h

@@ -16,76 +16,15 @@
 #define GPUSETTINGS_H
 
 #include "GPUCommonDef.h"
-#include "GPUDataTypesConfig.h"
-#include "GPUTPCGMMergedTrackHit.h"
-#ifndef GPUCA_GPUCODE_DEVICE
-#include <vector>
-#include <string>
-#endif
-
-namespace o2::gpu
-{
-class GPUDisplayFrontendInterface;
-class GPUReconstruction;
-
-class GPUSettings
-{
- public:
-  enum CompressionModes { CompressionTruncate = 1,
-                          CompressionDifferences = 2,
-                          CompressionTrackModel = 4,
-                          CompressionFull = 7 };
-  enum CompressionSort { SortTime = 0,
-                         SortPad = 1,
-                         SortZTimePad = 2,
-                         SortZPadTime = 3,
-                         SortNoSort = 4 };
-  enum CompressionRejection { RejectionNone = 0,
-                              RejectionStrategyA = 1,
-                              RejectionStrategyB = 2 };
-
-  static GPUglobalconstexpr() const uint32_t TPC_MAX_TF_TIME_BIN = ((256 * 3564 + 2 * 8 - 2) / 8);
-};
-
-// Settings describing the global run parameters
-struct GPUSettingsGRP {
-  // All new members must be sizeof(int32_t) resp. sizeof(float) for alignment reasons!, default value for newly added members for old data will be 0.
-  float solenoidBzNominalGPU = -5.00668f; // solenoid field strength
-  int32_t constBz = 0;                    // for test-MC events with constant Bz
-  int32_t removed0 = 0;                   // Obsolete parameter, dummy value needed to support reading old dumps
-  int32_t grpContinuousMaxTimeBin = -2;   // 0 for triggered events, -1 for automatic setting, -2 invalid default
-  int32_t needsClusterer = 0;             // Set to true if the data requires the clusterizer
-  int32_t doCompClusterDecode = 0;        // Set to true if the data contains compressed TPC clusters
-  int32_t tpcCutTimeBin = 0;              // Cut TPC clusters and digits >= this cut
-};
+#include "GPUSettingsTypes.h"
 
-// Parameters of the current time frame
-struct GPUSettingsTF {
-  int32_t hasTfStartOrbit = 0;
-  int32_t tfStartOrbit = 0;
-  int32_t hasRunStartOrbit = 0;
-  int32_t runStartOrbit = 0;
-  int32_t hasSimStartOrbit = 0;
-  int32_t simStartOrbit = 0;
-  int32_t hasNHBFPerTF = 0;
-  int32_t nHBFPerTF = 0;
-};
-
-// Settings defining the setup of the GPUReconstruction processing (basically selecting the device / class instance)
-struct GPUSettingsDeviceBackend {
-  uint32_t deviceType = gpudatatypes::DeviceType::CPU; // Device type, shall use gpudatatypes::DEVICE_TYPE constants, e.g. CPU / CUDA
-  uint8_t forceDeviceType = 1;                         // Fail if device initialization fails, otherwise falls back to CPU
-  GPUReconstruction* master = nullptr;                 // GPUReconstruction master object
-};
-
-} // namespace o2::gpu
+#endif // GPUSETTINGS_H
 
+// Outside of header guard, since might run it multiple times for different config structs
 #ifdef GPUCA_GPUCODE_DEVICE
 #define QCONFIG_GPU
 #endif
 // See GPUSettingsList.h for a list of all available settings of GPU Reconstruction
 #ifndef GPUCA_GPUCODE_GENRTC
 #include "utils/qconfig.h"
 #endif
-
-#endif

GPU/GPUTracking/DataTypes/GPUSettingsParam.h

@@ -0,0 +1,23 @@
+// 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 GPUSettingsParam.h
+/// \author David Rohr
+
+// clang-format off
+
+// No include guard: included multiple times by qconfig with different macro definitions
+
+// Derrived parameters used in GPUParam
+BeginHiddenConfig(GPUSettingsParam, param)
+AddVariableRTC(continuousTracking, int8_t, 0) // Continuous tracking, estimate bz and errors for abs(z) = 125cm during seeding
+AddVariableRTC(dodEdx, int8_t, 0)             // Do dEdx computation
+EndConfig()

GPU/GPUTracking/DataTypes/GPUSettingsTypes.h

@@ -0,0 +1,83 @@
+// 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 GPUSettingsTypes.h
+/// \author David Rohr
+
+#ifndef GPUSETTINGSTYPES_H
+#define GPUSETTINGSTYPES_H
+
+#include "GPUCommonDef.h"
+#include "GPUDataTypesConfig.h"
+#include "GPUTPCGMMergedTrackHit.h"
+#ifndef GPUCA_GPUCODE_DEVICE
+#include <vector>
+#include <string>
+#endif
+
+namespace o2::gpu
+{
+class GPUDisplayFrontendInterface;
+class GPUReconstruction;
+
+class GPUSettings
+{
+ public:
+  enum CompressionModes { CompressionTruncate = 1,
+                          CompressionDifferences = 2,
+                          CompressionTrackModel = 4,
+                          CompressionFull = 7 };
+  enum CompressionSort { SortTime = 0,
+                         SortPad = 1,
+                         SortZTimePad = 2,
+                         SortZPadTime = 3,
+                         SortNoSort = 4 };
+  enum CompressionRejection { RejectionNone = 0,
+                              RejectionStrategyA = 1,
+                              RejectionStrategyB = 2 };
+
+  static GPUglobalconstexpr() const uint32_t TPC_MAX_TF_TIME_BIN = ((256 * 3564 + 2 * 8 - 2) / 8);
+};
+
+// Settings describing the global run parameters
+struct GPUSettingsGRP {
+  // All new members must be sizeof(int32_t) resp. sizeof(float) for alignment reasons!, default value for newly added members for old data will be 0.
+  float solenoidBzNominalGPU = -5.00668f; // solenoid field strength
+  int32_t constBz = 0;                    // for test-MC events with constant Bz
+  int32_t removed0 = 0;                   // Obsolete parameter, dummy value needed to support reading old dumps
+  int32_t grpContinuousMaxTimeBin = -2;   // 0 for triggered events, -1 for automatic setting, -2 invalid default
+  int32_t needsClusterer = 0;             // Set to true if the data requires the clusterizer
+  int32_t doCompClusterDecode = 0;        // Set to true if the data contains compressed TPC clusters
+  int32_t tpcCutTimeBin = 0;              // Cut TPC clusters and digits >= this cut
+};
+
+// Parameters of the current time frame
+struct GPUSettingsTF {
+  int32_t hasTfStartOrbit = 0;
+  int32_t tfStartOrbit = 0;
+  int32_t hasRunStartOrbit = 0;
+  int32_t runStartOrbit = 0;
+  int32_t hasSimStartOrbit = 0;
+  int32_t simStartOrbit = 0;
+  int32_t hasNHBFPerTF = 0;
+  int32_t nHBFPerTF = 0;
+};
+
+// Settings defining the setup of the GPUReconstruction processing (basically selecting the device / class instance)
+struct GPUSettingsDeviceBackend {
+  uint32_t deviceType = gpudatatypes::DeviceType::CPU; // Device type, shall use gpudatatypes::DEVICE_TYPE constants, e.g. CPU / CUDA
+  uint8_t forceDeviceType = 1;                         // Fail if device initialization fails, otherwise falls back to CPU
+  GPUReconstruction* master = nullptr;                 // GPUReconstruction master object
+};
+
+} // namespace o2::gpu
+
+#endif // GPUSETTINGSTYPES_H