GPU: Better hiding of internal structures

Author: davidrohr

GPU/GPUTracking/Base/GPUReconstruction.cxx

@@ -46,6 +46,8 @@
 
 namespace o2::gpu
 {
+namespace // anonymous
+{
 struct GPUReconstructionPipelineQueue {
   uint32_t op = 0; // For now, 0 = process, 1 = terminate
   GPUChain* chain = nullptr;
@@ -54,6 +56,7 @@ struct GPUReconstructionPipelineQueue {
   bool done = false;
   int32_t retVal = 0;
 };
+} // namespace
 
 struct GPUReconstructionPipelineContext {
   std::queue<GPUReconstructionPipelineQueue*> queue;

GPU/GPUTracking/Base/GPUReconstructionConvert.cxx

@@ -174,6 +174,8 @@ int32_t GPUReconstructionConvert::GetMaxTimeBin(const GPUTrackingInOutZS& zspage
 // ------------------------------------------------- TPC ZS -------------------------------------------------
 
 #ifdef GPUCA_TPC_GEOMETRY_O2
+namespace o2::gpu
+{
 namespace // anonymous
 {
 
@@ -1306,19 +1308,9 @@ size_t zsEncoderRun<T>::compare(std::vector<zsPage>* buffer, std::vector<o2::tpc
 }
 
 } // anonymous namespace
+} // namespace o2::gpu
 #endif // GPUCA_TPC_GEOMETRY_O2
 
-namespace o2::gpu::internal
-{
-struct tmpReductionResult {
-  uint32_t totalPages = 0;
-  size_t totalSize = 0;
-  size_t nErrors = 0;
-  size_t digitsInput = 0;
-  size_t digitsEncoded = 0;
-};
-} // namespace o2::gpu::internal
-
 template <class S>
 void GPUReconstructionConvert::RunZSEncoder(const S& in, std::unique_ptr<uint64_t[]>* outBuffer, uint32_t* outSizes, o2::raw::RawFileWriter* raw, const o2::InteractionRecord* ir, const GPUParam& param, int32_t version, bool verify, float threshold, bool padding, std::function<void(std::vector<o2::tpc::Digit>&)> digitsFilter)
 {
@@ -1329,7 +1321,14 @@ void GPUReconstructionConvert::RunZSEncoder(const S& in, std::unique_ptr<uint64_
   }
 #ifdef GPUCA_TPC_GEOMETRY_O2
   std::vector<zsPage> buffer[NSECTORS][GPUTrackingInOutZS::NENDPOINTS];
-  auto reduced = tbb::parallel_reduce(tbb::blocked_range<uint32_t>(0, NSECTORS), o2::gpu::internal::tmpReductionResult(), [&](const auto range, auto red) {
+  struct tmpReductionResult {
+    uint32_t totalPages = 0;
+    size_t totalSize = 0;
+    size_t nErrors = 0;
+    size_t digitsInput = 0;
+    size_t digitsEncoded = 0;
+  };
+  auto reduced = tbb::parallel_reduce(tbb::blocked_range<uint32_t>(0, NSECTORS), tmpReductionResult(), [&](const auto range, auto red) {
     for (uint32_t i = range.begin(); i < range.end(); i++) {
       std::vector<o2::tpc::Digit> tmpBuffer;
       red.digitsInput += ZSEncoderGetNDigits(in, i);
@@ -1455,6 +1454,8 @@ void GPUReconstructionConvert::RunZSFilter(std::unique_ptr<o2::tpc::Digit[]>* bu
 }
 
 #ifdef GPUCA_O2_LIB
+namespace o2::gpu::internal
+{
 template <class T>
 static inline auto GetDecoder_internal(const GPUParam* param, int32_t version)
 {
@@ -1480,15 +1481,16 @@ static inline auto GetDecoder_internal(const GPUParam* param, int32_t version)
     enc->decodePage(outBuffer, (const zsPage*)page, endpoint, firstTfOrbit, triggerBC);
   };
 }
+} // namespace o2::gpu::internal
 
 std::function<void(std::vector<o2::tpc::Digit>&, const void*, uint32_t, uint32_t)> GPUReconstructionConvert::GetDecoder(int32_t version, const GPUParam* param)
 {
   if (version >= o2::tpc::ZSVersion::ZSVersionRowBased10BitADC && version <= o2::tpc::ZSVersion::ZSVersionRowBased12BitADC) {
-    return GetDecoder_internal<zsEncoderRow>(param, version);
+    return o2::gpu::internal::GetDecoder_internal<zsEncoderRow>(param, version);
   } else if (version == o2::tpc::ZSVersion::ZSVersionLinkBasedWithMeta) {
-    return GetDecoder_internal<zsEncoderImprovedLinkBased>(param, version);
+    return o2::gpu::internal::GetDecoder_internal<zsEncoderImprovedLinkBased>(param, version);
   } else if (version >= o2::tpc::ZSVersion::ZSVersionDenseLinkBased && version <= o2::tpc::ZSVersion::ZSVersionDenseLinkBasedV2) {
-    return GetDecoder_internal<zsEncoderDenseLinkBased>(param, version);
+    return o2::gpu::internal::GetDecoder_internal<zsEncoderDenseLinkBased>(param, version);
   } else {
     throw std::runtime_error("Invalid ZS version "s + std::to_string(version) + ", cannot create decoder"s);
   }

GPU/GPUTracking/Base/GPUReconstructionProcessing.cxx

@@ -50,7 +50,13 @@ void GPUReconstructionProcessing::runParallelOuterLoop(bool doGPU, uint32_t nThr
   });
 }
 
+namespace o2::gpu
+{
+namespace // anonymous
+{
 static std::atomic_flag timerFlag = ATOMIC_FLAG_INIT; // TODO: Should be a class member not global, but cannot be moved to header due to ROOT limitation
+} // anonymous namespace
+} // namespace o2::gpu
 
 GPUReconstructionProcessing::timerMeta* GPUReconstructionProcessing::insertTimer(uint32_t id, std::string&& name, int32_t J, int32_t num, int32_t type, RecoStep step)
 {

GPU/GPUTracking/DataCompression/GPUTPCClusterStatistics.cxx

@@ -23,7 +23,9 @@
 using namespace o2::gpu;
 
 // Small helper to compute Huffman probabilities
-namespace
+namespace o2::gpu
+{
+namespace // anonymous
 {
 typedef std::vector<bool> HuffCode;
 typedef std::map<uint32_t, HuffCode> HuffCodeMap;
@@ -101,7 +103,8 @@ void GenerateCodes(const INode* node, const HuffCode& prefix, HuffCodeMap& outCo
     GenerateCodes(in->right, rightPrefix, outCodes);
   }
 }
-} // namespace
+} // anonymous namespace
+} // namespace o2::gpu
 
 void GPUTPCClusterStatistics::RunStatistics(const o2::tpc::ClusterNativeAccess* clustersNative, const o2::tpc::CompressedClusters* clustersCompressed, const GPUParam& param)
 {

GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx

@@ -58,22 +58,26 @@
 #include "SimulationDataFormat/MCCompLabel.h"
 #endif
 
+namespace o2::gpu::internal
+{
+}
 using namespace o2::gpu;
+using namespace o2::gpu::internal;
 using namespace o2::tpc;
 using namespace gputpcgmmergertypes;
 
 static constexpr int32_t kMaxParts = 400;
 static constexpr int32_t kMaxClusters = GPUCA_MERGER_MAX_TRACK_CLUSTERS;
 
-namespace o2::gpu
+namespace o2::gpu::internal
 {
 struct MergeLooperParam {
   float refz;
   float x;
   float y;
   uint32_t id;
 };
-} // namespace o2::gpu
+} // namespace o2::gpu::internal
 
 #ifndef GPUCA_GPUCODE
 
@@ -741,6 +745,10 @@ GPUd() void GPUTPCGMMerger::MergeBorderTracks<1>(int32_t nBlocks, int32_t nThrea
 }
 
 #if defined(GPUCA_SPECIALIZE_THRUST_SORTS) && !defined(GPUCA_GPUCODE_COMPILEKERNELS) // Specialize MergeBorderTracks<3>
+namespace o2::gpu::internal
+{
+namespace // anonymous
+{
 struct MergeBorderTracks_compMax {
   GPUd() bool operator()(const GPUTPCGMBorderRange& a, const GPUTPCGMBorderRange& b)
   {
@@ -761,6 +769,8 @@ struct MergeBorderTracks_compMin {
 #endif
   }
 };
+} // anonymous namespace
+} // namespace o2::gpu::internal
 
 template <>
 inline void GPUCA_KRNL_BACKEND_CLASS::runKernelBackendInternal<GPUTPCGMMergerMergeBorders, 3>(const krnlSetupTime& _xyz, GPUTPCGMBorderRange* const& range, int32_t const& N, int32_t const& cmpMax)
@@ -1436,6 +1446,10 @@ GPUd() void GPUTPCGMMerger::MergeCE(int32_t nBlocks, int32_t nThreads, int32_t i
   // for (int32_t i = 0;i < mMemory->nOutputTracks;i++) {if (mOutputTracks[i].CCE() == false) {mOutputTracks[i].SetNClusters(0);mOutputTracks[i].SetOK(false);}} //Remove all non-CE tracks
 }
 
+namespace o2::gpu::internal
+{
+namespace // anonymous
+{
 struct GPUTPCGMMerger_CompareClusterIdsLooper {
   struct clcomparestruct {
     uint8_t leg;
@@ -1489,6 +1503,8 @@ struct GPUTPCGMMerger_CompareClusterIds {
 #endif
   }
 };
+} // anonymous namespace
+} // namespace o2::gpu::internal
 
 GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
 {
@@ -1803,6 +1819,10 @@ GPUd() void GPUTPCGMMerger::PrepareClustersForFit0(int32_t nBlocks, int32_t nThr
 }
 
 #if defined(GPUCA_SPECIALIZE_THRUST_SORTS) && !defined(GPUCA_GPUCODE_COMPILEKERNELS) // Specialize GPUTPCGMMergerSortTracks and GPUTPCGMMergerSortTracksQPt
+namespace o2::gpu::internal
+{
+namespace // anonymous
+{
 struct GPUTPCGMMergerSortTracks_comp {
   const GPUTPCGMMergedTrack* const mCmp;
   GPUhd() GPUTPCGMMergerSortTracks_comp(GPUTPCGMMergedTrack* cmp) : mCmp(cmp) {}
@@ -1833,14 +1853,6 @@ struct GPUTPCGMMergerSortTracks_comp {
   }
 };
 
-template <>
-inline void GPUCA_KRNL_BACKEND_CLASS::runKernelBackendInternal<GPUTPCGMMergerSortTracks, 0>(const krnlSetupTime& _xyz)
-{
-  thrust::device_ptr<uint32_t> trackSort((uint32_t*)mProcessorsShadow->tpcMerger.TrackOrderProcess());
-  ThrustVolatileAsyncAllocator alloc(this);
-  thrust::sort(GPUCA_THRUST_NAMESPACE::par(alloc).on(mInternals->Streams[_xyz.x.stream]), trackSort, trackSort + processors()->tpcMerger.NOutputTracks(), GPUTPCGMMergerSortTracks_comp(mProcessorsShadow->tpcMerger.OutputTracks()));
-}
-
 struct GPUTPCGMMergerSortTracksQPt_comp {
   const GPUTPCGMMergedTrack* const mCmp;
   GPUhd() GPUTPCGMMergerSortTracksQPt_comp(GPUTPCGMMergedTrack* cmp) : mCmp(cmp) {}
@@ -1861,6 +1873,16 @@ struct GPUTPCGMMergerSortTracksQPt_comp {
 #endif
   }
 };
+} // anonymous namespace
+} // namespace o2::gpu::internal
+
+template <>
+inline void GPUCA_KRNL_BACKEND_CLASS::runKernelBackendInternal<GPUTPCGMMergerSortTracks, 0>(const krnlSetupTime& _xyz)
+{
+  thrust::device_ptr<uint32_t> trackSort((uint32_t*)mProcessorsShadow->tpcMerger.TrackOrderProcess());
+  ThrustVolatileAsyncAllocator alloc(this);
+  thrust::sort(GPUCA_THRUST_NAMESPACE::par(alloc).on(mInternals->Streams[_xyz.x.stream]), trackSort, trackSort + processors()->tpcMerger.NOutputTracks(), GPUTPCGMMergerSortTracks_comp(mProcessorsShadow->tpcMerger.OutputTracks()));
+}
 
 template <>
 inline void GPUCA_KRNL_BACKEND_CLASS::runKernelBackendInternal<GPUTPCGMMergerSortTracksQPt, 0>(const krnlSetupTime& _xyz)
@@ -2074,12 +2096,18 @@ GPUd() void GPUTPCGMMerger::MergeLoopersSort(int32_t nBlocks, int32_t nThreads,
 }
 
 #if defined(GPUCA_SPECIALIZE_THRUST_SORTS) && !defined(GPUCA_GPUCODE_COMPILEKERNELS) // Specialize GPUTPCGMMergerSortTracks and GPUTPCGMMergerSortTracksQPt
+namespace o2::gpu::internal
+{
+namespace // anonymous
+{
 struct GPUTPCGMMergerMergeLoopers_comp {
   GPUd() bool operator()(const MergeLooperParam& a, const MergeLooperParam& b)
   {
     return CAMath::Abs(a.refz) < CAMath::Abs(b.refz);
   }
 };
+} // anonymous namespace
+} // namespace o2::gpu::internal
 
 template <>
 inline void GPUCA_KRNL_BACKEND_CLASS::runKernelBackendInternal<GPUTPCGMMergerMergeLoopers, 1>(const krnlSetupTime& _xyz)

GPU/GPUTracking/Merger/GPUTPCGMMerger.h

@@ -48,7 +48,10 @@ class GPUTPCTracker;
 class GPUChainTracking;
 class GPUTPCGMPolynomialField;
 struct GPUTPCGMLoopData;
+namespace internal
+{
 struct MergeLooperParam;
+} // namespace internal
 
 /**
  * @class GPUTPCGMMerger
@@ -126,7 +129,7 @@ class GPUTPCGMMerger : public GPUProcessor
   GPUhdi() uint2* ClusRefTmp() { return mClusRefTmp; }
   GPUhdi() uint32_t* TrackSort() { return mTrackSort; }
   GPUhdi() tmpSort* TrackSortO2() { return mTrackSortO2; }
-  GPUhdi() MergeLooperParam* LooperCandidates() { return mLooperCandidates; }
+  GPUhdi() internal::MergeLooperParam* LooperCandidates() { return mLooperCandidates; }
   GPUhdi() GPUAtomic(uint32_t) * SharedCount() { return mSharedCount; }
   GPUhdi() gputpcgmmergertypes::GPUTPCGMBorderRange* BorderRange(int32_t i) { return mBorderRange[i]; }
   GPUhdi() const gputpcgmmergertypes::GPUTPCGMBorderRange* BorderRange(int32_t i) const { return mBorderRange[i]; }
@@ -267,7 +270,7 @@ class GPUTPCGMMerger : public GPUProcessor
   o2::tpc::TrackTPC* mOutputTracksTPCO2;
   uint32_t* mOutputClusRefsTPCO2;
   o2::MCCompLabel* mOutputTracksTPCO2MC;
-  MergeLooperParam* mLooperCandidates;
+  internal::MergeLooperParam* mLooperCandidates;
 
   uint32_t* mTrackOrderAttach;
   uint32_t* mTrackOrderProcess;

GPU/GPUTracking/Refit/GPUTrackingRefit.cxx

@@ -63,7 +63,9 @@ void GPUTrackingRefitProcessor::SetMaxData(const GPUTrackingInOutPointers& io)
 }
 #endif
 
-namespace
+namespace o2::gpu::internal
+{
+namespace // anonymous
 {
 template <class T>
 struct refitTrackTypes;
@@ -76,6 +78,7 @@ struct refitTrackTypes<TrackParCov> {
   using propagator = const Propagator*;
 };
 } // anonymous namespace
+} // namespace o2::gpu::internal
 
 template <>
 GPUd() void GPUTrackingRefit::initProp<GPUgeneric() GPUTPCGMPropagator>(GPUTPCGMPropagator& prop) // FIXME: GPUgeneric() needed to make the clang spirv output link correctly
@@ -210,10 +213,10 @@ template <class T, class S>
 GPUd() int32_t GPUTrackingRefit::RefitTrack(T& trkX, bool outward, bool resetCov)
 {
   CADEBUG(int32_t ii; printf("\nRefitting track\n"));
-  typename refitTrackTypes<S>::propagator prop;
+  typename internal::refitTrackTypes<S>::propagator prop;
   S trk;
   float TrackParCovChi2 = 0.f;
-  convertTrack<S, T, typename refitTrackTypes<S>::propagator>(trk, trkX, prop, &TrackParCovChi2);
+  convertTrack<S, T, typename internal::refitTrackTypes<S>::propagator>(trk, trkX, prop, &TrackParCovChi2);
   int32_t begin = 0, count;
   float tOffset;
   if constexpr (std::is_same_v<T, GPUTPCGMMergedTrack>) {
@@ -417,7 +420,7 @@ GPUd() int32_t GPUTrackingRefit::RefitTrack(T& trkX, bool outward, bool resetCov
     static_assert("Invalid template");
   }
 
-  convertTrack<T, S, typename refitTrackTypes<S>::propagator>(trkX, trk, prop, &TrackParCovChi2);
+  convertTrack<T, S, typename internal::refitTrackTypes<S>::propagator>(trkX, trk, prop, &TrackParCovChi2);
   return nFitted;
 }
 

GPU/GPUTracking/display/frontend/GPUDisplayFrontendWayland.cxx

@@ -39,16 +39,8 @@
 
 using namespace o2::gpu;
 
-GPUDisplayFrontendWayland::GPUDisplayFrontendWayland()
-{
-  mFrontendType = TYPE_WAYLAND;
-  mFrontendName = "Wayland";
-}
-
-void GPUDisplayFrontendWayland::OpenGLPrint(const char* s, float x, float y, float r, float g, float b, float a, bool fromBotton)
+namespace o2::gpu::internal
 {
-}
-
 template <class T, class... Args>
 struct CCallWrapper {
   std::function<T(Args...)> func;
@@ -58,6 +50,17 @@ struct CCallWrapper {
     return funcwrap->func(std::forward<Args>(args)...);
   }
 };
+} // namespace o2::gpu::internal
+
+GPUDisplayFrontendWayland::GPUDisplayFrontendWayland()
+{
+  mFrontendType = TYPE_WAYLAND;
+  mFrontendName = "Wayland";
+}
+
+void GPUDisplayFrontendWayland::OpenGLPrint(const char* s, float x, float y, float r, float g, float b, float a, bool fromBotton)
+{
+}
 
 int32_t GPUDisplayFrontendWayland::GetKey(uint32_t key, uint32_t state)
 {
@@ -283,7 +286,7 @@ int32_t GPUDisplayFrontendWayland::FrontendMain()
       wl_keyboard_add_listener(mKeyboard, &keyboard_listener, this);
     }
   };
-  auto seat_capabilities_c = CCallWrapper<void, wl_seat*, uint32_t>{[seat_capabilities](wl_seat* seat, uint32_t capabilities) { seat_capabilities(seat, capabilities); }};
+  auto seat_capabilities_c = internal::CCallWrapper<void, wl_seat*, uint32_t>{[seat_capabilities](wl_seat* seat, uint32_t capabilities) { seat_capabilities(seat, capabilities); }};
 
   auto seat_name = [](void* data, struct wl_seat* seat, const char* name) {
     if (((GPUDisplayFrontendWayland*)data)->mDisplay->param()->par.debugLevel >= 2) {
@@ -317,7 +320,7 @@ int32_t GPUDisplayFrontendWayland::FrontendMain()
     }
   };
 
-  auto registry_global_c = CCallWrapper<void, wl_registry*, uint32_t, const char*, uint32_t>{[registry_global](wl_registry* registry, uint32_t name, const char* interface, uint32_t version) { registry_global(registry, name, interface, version); }};
+  auto registry_global_c = internal::CCallWrapper<void, wl_registry*, uint32_t, const char*, uint32_t>{[registry_global](wl_registry* registry, uint32_t name, const char* interface, uint32_t version) { registry_global(registry, name, interface, version); }};
   auto registry_global_remove = [](void* a, wl_registry* b, uint32_t c) {};
   const wl_registry_listener registry_listener = {.global = &registry_global_c.callback, .global_remove = registry_global_remove};