GPU: Clean up code that is no longer used

DataFormats/Detectors/TPC/include/DataFormatsTPC/CalibdEdxCorrection.h

@@ -49,9 +49,9 @@ class CalibdEdxCorrection
   }
   CalibdEdxCorrection(std::string_view fileName) { loadFromFile(fileName); }
 #else
-  CalibdEdxCorrection() CON_DEFAULT;
+  CalibdEdxCorrection() = default;
 #endif
-  ~CalibdEdxCorrection() CON_DEFAULT;
+  ~CalibdEdxCorrection() = default;
 
   GPUd() float getCorrection(const StackID& stack, ChargeType charge, float tgl = 0, float snp = 0) const
   {

DataFormats/Detectors/TPC/include/DataFormatsTPC/ClusterNative.h

@@ -76,7 +76,7 @@ struct ClusterNative {
   GPUd() static float unpackPad(uint16_t pad) { return float(pad) * (1.f / scalePadPacked); }
   GPUd() static float unpackTime(uint32_t time) { return float(time) * (1.f / scaleTimePacked); }
 
-  GPUdDefault() ClusterNative() CON_DEFAULT;
+  GPUdDefault() ClusterNative() = default;
   GPUd() ClusterNative(uint32_t time, uint8_t flags, uint16_t pad, uint8_t sigmaTime, uint8_t sigmaPad, uint16_t qmax, uint16_t qtot) : padPacked(pad), sigmaTimePacked(sigmaTime), sigmaPadPacked(sigmaPad), qMax(qmax), qTot(qtot)
   {
     setTimePackedFlags(time, flags);

DataFormats/Detectors/TPC/include/DataFormatsTPC/CompressedClusters.h

@@ -77,8 +77,8 @@ struct CompressedClustersOffsets : public CompressedClustersPtrs_x<size_t, size_
 struct CompressedClustersFlat;
 
 struct CompressedClusters : public CompressedClustersCounters, public CompressedClustersPtrs { // TODO: Need a const version of this, currently the constructor allows to create a non-const version from const CompressedClustersFlat, which should not be allowed
-  CompressedClusters() CON_DEFAULT;
-  ~CompressedClusters() CON_DEFAULT;
+  CompressedClusters() = default;
+  ~CompressedClusters() = default;
   CompressedClusters(const CompressedClustersFlat& c);
 
   void dump();
@@ -87,7 +87,7 @@ struct CompressedClusters : public CompressedClustersCounters, public Compressed
 };
 
 struct CompressedClustersROOT : public CompressedClusters {
-  CompressedClustersROOT() CON_DEFAULT;
+  CompressedClustersROOT() = default;
   CompressedClustersROOT(const CompressedClustersFlat& v) : CompressedClusters(v) {}
   CompressedClustersROOT(const CompressedClusters& v) : CompressedClusters(v) {}
   // flatbuffer used for streaming
@@ -99,7 +99,7 @@ struct CompressedClustersROOT : public CompressedClusters {
 
 struct CompressedClustersFlat : private CompressedClustersCounters, private CompressedClustersOffsets {
   friend struct CompressedClusters;               // We don't want anyone to access the members directly, should only be used to construct a CompressedClusters struct
-  CompressedClustersFlat() CON_DELETE;            // Must not be constructed
+  CompressedClustersFlat() = delete;              // Must not be constructed
   size_t totalDataSize = 0;                       // Total data size of header + content
   const CompressedClusters* ptrForward = nullptr; // Must be 0 if this object is really flat, or can be a ptr to a CompressedClusters struct (abusing the flat structure to forward a ptr to the e.g. root version)
 

DataFormats/Detectors/TPC/include/DataFormatsTPC/Digit.h

@@ -36,7 +36,7 @@ class Digit : public DigitBase
 {
  public:
   /// Default constructor
-  GPUdDefault() Digit() CON_DEFAULT;
+  GPUdDefault() Digit() = default;
 
   /// Constructor, initializing values for position, charge, time and common mode
   /// \param cru CRU of the Digit
@@ -46,7 +46,7 @@ class Digit : public DigitBase
   GPUdi() Digit(int cru, float charge, int row, int pad, int time);
 
   /// Destructor
-  GPUdDefault() ~Digit() CON_DEFAULT;
+  GPUdDefault() ~Digit() = default;
 
   /// Get the accumulated charged of the Digit in ADC counts.
   /// The conversion is such that the decimals are simply stripped

DataFormats/Detectors/TPC/include/DataFormatsTPC/PIDResponse.h

@@ -39,10 +39,10 @@ class PIDResponse
 {
  public:
   /// default constructor
-  PIDResponse() CON_DEFAULT;
+  PIDResponse() = default;
 
   /// default destructor
-  ~PIDResponse() CON_DEFAULT;
+  ~PIDResponse() = default;
 
   /// setters
   GPUd() void setBetheBlochParams(const float betheBlochParams[5]);

DataFormats/common/include/CommonDataFormat/TimeStamp.h

@@ -25,8 +25,8 @@ template <typename T>
 class TimeStamp
 {
  public:
-  GPUhdDefault() TimeStamp() CON_DEFAULT;
-  GPUhdDefault() ~TimeStamp() CON_DEFAULT;
+  GPUhdDefault() TimeStamp() = default;
+  GPUhdDefault() ~TimeStamp() = default;
   GPUdi() TimeStamp(T time) { mTimeStamp = time; }
   GPUhdi() T getTimeStamp() const { return mTimeStamp; }
   GPUdi() void setTimeStamp(T t) { mTimeStamp = t; }

Detectors/Base/include/DetectorsBase/MatCell.h

@@ -31,7 +31,7 @@ struct MatCell {
   float meanX2X0;                   ///< fraction of radiaton lenght
 
   GPUd() MatCell() : meanRho(0.f), meanX2X0(0.f) {}
-  GPUdDefault() MatCell(const MatCell& src) CON_DEFAULT;
+  GPUdDefault() MatCell(const MatCell& src) = default;
 
   GPUd() void set(const MatCell& c)
   {
@@ -55,7 +55,7 @@ struct MatBudget : MatCell {
   float length;                     ///< length in material
 
   GPUd() MatBudget() : length(0.f) {}
-  GPUdDefault() MatBudget(const MatBudget& src) CON_DEFAULT;
+  GPUdDefault() MatBudget(const MatBudget& src) = default;
 
   GPUd() void scale(float scale)
   {

Detectors/Base/include/DetectorsBase/MatLayerCyl.h

@@ -56,8 +56,8 @@ class MatLayerCyl : public o2::gpu::FlatObject
 
 #ifndef GPUCA_GPUCODE
   MatLayerCyl();
-  MatLayerCyl(const MatLayerCyl& src) CON_DELETE;
-  ~MatLayerCyl() CON_DEFAULT;
+  MatLayerCyl(const MatLayerCyl& src) = delete;
+  ~MatLayerCyl() = default;
 #endif
 
 #ifndef GPUCA_ALIGPUCODE // this part is unvisible on GPU version

Detectors/Base/include/DetectorsBase/MatLayerCylSet.h

@@ -52,9 +52,9 @@ class MatLayerCylSet : public o2::gpu::FlatObject
 
  public:
 #ifndef GPUCA_GPUCODE
-  MatLayerCylSet() CON_DEFAULT;
-  ~MatLayerCylSet() CON_DEFAULT;
-  MatLayerCylSet(const MatLayerCylSet& src) CON_DELETE;
+  MatLayerCylSet() = default;
+  ~MatLayerCylSet() = default;
+  MatLayerCylSet(const MatLayerCylSet& src) = delete;
 #endif
 
   GPUd() const MatLayerCylSetLayout* get() const { return reinterpret_cast<const MatLayerCylSetLayout*>(mFlatBufferPtr); }

Detectors/Base/include/DetectorsBase/Ray.h

@@ -49,7 +49,7 @@ class Ray
   GPUd() Ray() : mP{0.f}, mD{0.f}, mDistXY2(0.f), mDistXY2i(0.f), mDistXYZ(0.f), mXDxPlusYDy(0.f), mXDxPlusYDyRed(0.f), mXDxPlusYDy2(0.f), mR02(0.f), mR12(0.f)
   {
   }
-  GPUdDefault() ~Ray() CON_DEFAULT;
+  GPUdDefault() ~Ray() = default;
 
 #ifndef GPUCA_ALIGPUCODE // this part is unvisible on GPU version
   Ray(const math_utils::Point3D<float> point0, const math_utils::Point3D<float> point1);

Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu

@@ -1215,14 +1215,14 @@ void processNeighboursHandler(const int startLayer,
                                               thrust::raw_pointer_cast(&foundSeedsTable[0]), // d_in
                                               thrust::raw_pointer_cast(&foundSeedsTable[0]), // d_out
                                               nCurrentCells + 1,                             // num_items
-                                              0));
+                                              0));                                           // NOLINT: failure in clang-tidy
   discardResult(cudaMalloc(&d_temp_storage, temp_storage_bytes));
   gpuCheckError(cub::DeviceScan::ExclusiveSum(d_temp_storage,                                // d_temp_storage
                                               temp_storage_bytes,                            // temp_storage_bytes
                                               thrust::raw_pointer_cast(&foundSeedsTable[0]), // d_in
                                               thrust::raw_pointer_cast(&foundSeedsTable[0]), // d_out
                                               nCurrentCells + 1,                             // num_items
-                                              0));
+                                              0));                                           // NOLINT: failure in clang-tidy
 
   thrust::device_vector<int> updatedCellIds(foundSeedsTable.back()) /*, lastCellIds(foundSeedsTable.back())*/;
   thrust::device_vector<CellSeed> updatedCellSeeds(foundSeedsTable.back()) /*, lastCellSeeds(foundSeedsTable.back())*/;

GPU/Common/GPUCommonConstants.h

@@ -17,11 +17,9 @@
 
 #include "GPUCommonDef.h"
 
-#if !defined(__OPENCL1__)
 namespace GPUCA_NAMESPACE::gpu::gpu_common_constants
 {
-static CONSTEXPR const float kCLight = 0.000299792458f; // TODO: Duplicate of MathConstants, fix this when OpenCL1 is removed
+static CONSTEXPR const float kCLight = 0.000299792458f; // TODO: Duplicate of MathConstants, fix this now that we use only OpenCL CPP
 }
-#endif
 
 #endif

GPU/Common/GPUCommonDef.h

@@ -30,17 +30,7 @@
 //Some GPU configuration settings, must be included first
 #include "GPUCommonDefSettings.h"
 
-#if !defined(__OPENCL1__) && (!(defined(__CINT__) || defined(__ROOTCINT__)) || defined(__CLING__)) && defined(__cplusplus) && __cplusplus >= 201103L
-  #define GPUCA_NOCOMPAT // C++11 + No old ROOT5 + No old OpenCL
-  #ifndef __OPENCL__
-    #define GPUCA_NOCOMPAT_ALLOPENCL // + No OpenCL at all
-  #endif
-  #ifndef __CINT__
-    #define GPUCA_NOCOMPAT_ALLCINT // + No ROOT CINT at all
-  #endif
-#endif
-
-#if !(defined(__CINT__) || defined(__ROOTCINT__) || defined(__CLING__) || defined(__ROOTCLING__) || defined(G__ROOT)) // No GPU code for ROOT
+#if !(defined(__CLING__) || defined(__ROOTCLING__) || defined(G__ROOT)) // No GPU code for ROOT
   #if defined(__CUDACC__) || defined(__OPENCL__) || defined(__HIPCC__) || defined(__OPENCL_HOST__)
     #define GPUCA_GPUCODE // Compiled by GPU compiler
   #endif
@@ -50,26 +40,11 @@
   #endif
 #endif
 
-// Definitions for C++11 features not supported by CINT / OpenCL
-#ifdef GPUCA_NOCOMPAT
-  #define CON_DELETE = delete
-  #define CON_DEFAULT = default
-  #define GPUCA_CPP11_INIT(...) __VA_ARGS__
-  #if defined(__cplusplus) && __cplusplus >= 201703L
-    #define CONSTEXPR constexpr
-  #else
-    #define CONSTEXPR
-  #endif
+// Definitions for C++11 features
+#if defined(__cplusplus) && __cplusplus >= 201703L
+  #define CONSTEXPR constexpr
 #else
-  #define CON_DELETE
-  #define CON_DEFAULT
   #define CONSTEXPR
-  #define GPUCA_CPP11_INIT(...)
-#endif
-#if defined(__ROOT__) && !defined(GPUCA_NOCOMPAT)
-  #define VOLATILE // ROOT5 has a problem with volatile in CINT
-#else
-  #define VOLATILE volatile
 #endif
 
 // Set AliRoot / O2 namespace
@@ -82,7 +57,7 @@
   #define GPUCA_NAMESPACE o2
 #endif
 
-#if (defined(__CUDACC__) && defined(GPUCA_CUDA_NO_CONSTANT_MEMORY)) || (defined(__HIPCC__) && defined(GPUCA_HIP_NO_CONSTANT_MEMORY)) || (defined(__OPENCL1__) && defined(GPUCA_OPENCL_NO_CONSTANT_MEMORY)) || (defined(__OPENCLCPP__) && defined(GPUCA_OPENCLCPP_NO_CONSTANT_MEMORY))
+#if (defined(__CUDACC__) && defined(GPUCA_CUDA_NO_CONSTANT_MEMORY)) || (defined(__HIPCC__) && defined(GPUCA_HIP_NO_CONSTANT_MEMORY)) || (defined(__OPENCL__) && defined(GPUCA_OPENCL_NO_CONSTANT_MEMORY))
   #define GPUCA_NO_CONSTANT_MEMORY
 #elif defined(__CUDACC__) || defined(__HIPCC__)
   #define GPUCA_HAS_GLOBAL_SYMBOL_CONSTANT_MEM

GPU/Common/GPUCommonDefAPI.h

@@ -95,15 +95,15 @@
   #define GPUprivate() __private
   #define GPUgeneric() __generic
   #define GPUconstexprref() GPUconstexpr()
-  #if defined(__OPENCLCPP__) && !defined(__clang__)
+  #if defined(__OPENCL__) && !defined(__clang__)
     #define GPUbarrier() work_group_barrier(mem_fence::global | mem_fence::local);
     #define GPUbarrierWarp()
     #define GPUAtomic(type) atomic<type>
     static_assert(sizeof(atomic<uint32_t>) == sizeof(uint32_t), "Invalid size of atomic type");
   #else
     #define GPUbarrier() barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE)
     #define GPUbarrierWarp()
-    #if defined(__OPENCLCPP__) && defined(GPUCA_OPENCL_CPP_CLANG_C11_ATOMICS)
+    #if defined(__OPENCL__) && defined(GPUCA_OPENCL_CLANG_C11_ATOMICS)
       namespace GPUCA_NAMESPACE { namespace gpu {
       template <class T> struct oclAtomic;
       template <> struct oclAtomic<uint32_t> {typedef atomic_uint t;};
@@ -114,14 +114,14 @@
       #define GPUAtomic(type) volatile type
     #endif
   #endif
-  #if !defined(__OPENCLCPP__) // Other special defines for OpenCL 1
+  #if !defined(__OPENCL__) // Other special defines for OpenCL 1
     #define GPUCA_USE_TEMPLATE_ADDRESS_SPACES // TODO: check if we can make this (partially, where it is already implemented) compatible with OpenCL CPP
     #define GPUsharedref() GPUshared()
     #define GPUglobalref() GPUglobal()
     #undef GPUgeneric
     #define GPUgeneric()
   #endif
-  #if (!defined(__OPENCLCPP__) || !defined(GPUCA_NO_CONSTANT_MEMORY))
+  #if (!defined(__OPENCL__) || !defined(GPUCA_NO_CONSTANT_MEMORY))
     #define GPUconstantref() GPUconstant()
   #endif
 #elif defined(__HIPCC__) //Defines for HIP

GPU/Common/GPUCommonDefSettings.h

@@ -22,12 +22,11 @@
   #error Please include GPUCommonDef.h!
 #endif
 
-//#define GPUCA_OPENCL_CPP_CLANG_C11_ATOMICS     // Use C11 atomic instead of old style atomics for OpenCL C++ in clang (OpenCL 2.2 C++ will use C++11 atomics irrespectively)
+//#define GPUCA_OPENCL_CLANG_C11_ATOMICS     // Use C11 atomic instead of old style atomics for OpenCL C++ in clang (OpenCL 2.2 C++ will use C++11 atomics irrespectively)
 
 //#define GPUCA_CUDA_NO_CONSTANT_MEMORY          // Do not use constant memory for CUDA
 //#define GPUCA_HIP_NO_CONSTANT_MEMORY           // Do not use constant memory for HIP
-//#define GPUCA_OPENCL_NO_CONSTANT_MEMORY        // Do not use constant memory for OpenCL 1.2
-#define GPUCA_OPENCLCPP_NO_CONSTANT_MEMORY       // Do not use constant memory for OpenCL C++ - MANDATORY as OpenCL cannot cast between __constant and __generic yet!
+#define GPUCA_OPENCL_NO_CONSTANT_MEMORY       // Do not use constant memory for OpenCL C++ - MANDATORY as OpenCL cannot cast between __constant and __generic yet!
 
 // clang-format on