Fix some HIP specific issues

Author: mconcas

DataFormats/Reconstruction/include/ReconstructionDataFormats/Vertex.h

@@ -40,7 +40,7 @@ class VertexBase
                         kCovYZ,
                         kCovZZ };
   static constexpr int kNCov = 6;
-  GPUdDefault() VertexBase() = default;
+  GPUhd() VertexBase() = default;
   GPUdDefault() ~VertexBase() = default;
   GPUd() VertexBase(const math_utils::Point3D<float>& pos, const gpu::gpustd::array<float, kNCov>& cov) : mPos(pos), mCov(cov)
   {
@@ -119,7 +119,7 @@ class Vertex : public VertexBase
     FlagsMask = 0xffff
   };
 
-  GPUdDefault() Vertex() = default;
+  GPUhd() Vertex() = default;
   GPUdDefault() ~Vertex() = default;
   GPUd() Vertex(const math_utils::Point3D<float>& pos, const gpu::gpustd::array<float, kNCov>& cov, ushort nCont, float chi2)
     : VertexBase(pos, cov), mNContributors(nCont), mChi2(chi2)

DataFormats/common/include/CommonDataFormat/TimeStamp.h

@@ -25,7 +25,7 @@ template <typename T>
 class TimeStamp
 {
  public:
-  GPUdDefault() TimeStamp() CON_DEFAULT;
+  GPUhd() TimeStamp() CON_DEFAULT;
   GPUdDefault() ~TimeStamp() CON_DEFAULT;
   GPUdi() TimeStamp(T time) { mTimeStamp = time; }
   GPUdi() T getTimeStamp() const { return mTimeStamp; }

Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h

@@ -73,7 +73,11 @@ class TimeFrameGPU : public TimeFrame
   int* getDeviceXHistograms(const int rofId);
   int* getDeviceYHistograms(const int rofId);
   int* getDeviceZHistograms(const int rofId);
+#ifdef __HIPCC__
+  hipcub::KeyValuePair<int, int>* getTmpVertexPositionBins(const int rofId);
+#else
   cub::KeyValuePair<int, int>* getTmpVertexPositionBins(const int rofId);
+#endif
   float* getDeviceBeamPosition(const int rofId);
   Vertex* getDeviceVertices(const int rofId);
 
@@ -100,9 +104,13 @@ class TimeFrameGPU : public TimeFrame
   Vector<float> mZCentroids;
   std::array<Vector<int>, 2> mNTrackletsPerClusterD;
   std::array<Vector<int>, 3> mXYZHistograms;
-  Vector<cub::KeyValuePair<int, int>> mTmpVertexPositionBins;
   Vector<float> mBeamPosition;
   Vector<Vertex> mGPUVertices;
+#ifdef __HIPCC__
+  Vector<hipcub::KeyValuePair<int, int>> mTmpVertexPositionBins;
+#else
+  Vector<cub::KeyValuePair<int, int>> mTmpVertexPositionBins;
+#endif
 };
 
 template <int NLayers>
@@ -235,7 +243,11 @@ inline int* TimeFrameGPU<NLayers>::getDeviceZHistograms(const int rofId)
 }
 
 template <int NLayers>
+#ifdef __HIPCC__
+inline hipcub::KeyValuePair<int, int>* TimeFrameGPU<NLayers>::getTmpVertexPositionBins(const int rofId)
+#else
 inline cub::KeyValuePair<int, int>* TimeFrameGPU<NLayers>::getTmpVertexPositionBins(const int rofId)
+#endif
 {
   if (rofId < 0 || rofId >= mNrof) {
     LOG(error) << "Invalid rofId: " << rofId << "/" << mNrof << ", returning nullptr";

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

@@ -342,7 +342,7 @@ GPUg() void computeVertexKernel(cub::KeyValuePair<int, int>* tmpVertexBins,
           histZ[iBin] = 0;
         }
         if (sumWZ > minContributors || vertIndex == 0) {
-          new (vertices + vertIndex) Vertex{o2::math_utils::Point3D<float>(beamPosition[0], beamPosition[1], wZ / sumWZ), o2::gpu::gpustd::array<float, 6>{ex, 0, ey, 0, 0, ez}, sumWZ, 0};
+          new (vertices + vertIndex) Vertex{o2::math_utils::Point3D<float>(beamPosition[0], beamPosition[1], wZ / sumWZ), o2::gpu::gpustd::array<float, 6>{ex, 0, ey, 0, 0, ez}, static_cast<ushort>(sumWZ), 0};
         } else {
           new (vertices + vertIndex) Vertex{};
         }