Integrate the comments from the upstream PRs (#442)

Clean up the Patatrack code base following the comments received during the integration into the upstream release.

Currently tracks the changes introduced due to
   - cms-sw#29109: Patatrack integration - trivial changes (1/N)
   - cms-sw#29110: Patatrack integration - common tools (2/N)

List of changes:
 * Remove unused files
 * Fix compilation warnings
 * Fix AtomicPairCounter unit test
 * Rename the cudaCompat namespace to cms::cudacompat
 * Remove extra semicolon
 * Move SimpleVector and VecArray to the cms::cuda namespace
 * Add missing dependency
 * Move HistoContainer, AtomicPairCounter, prefixScan and radixSort to the cms::cuda namespace
 * Remove rule exception for HeterogeneousCore
 * Fix code rule violations:
    - replace using namespace cms::cuda in test/OneToManyAssoc_t.h .
    - add an exception for cudaCompat.h:
      cudaCompat relies on defining equivalent symbols to the CUDA
      intrinsics in the cms::cudacompat namespace, and pulling them in the
      global namespace when compiling device code without CUDA.
* Protect the headers to compile only with a CUDA compiler

RecoPixelVertexing/PixelVertexFinding/plugins/gpuClusterTracksByDensity.h

@@ -48,7 +48,7 @@ namespace gpuVertexFinder {
     assert(pdata);
     assert(zt);
 
-    using Hist = HistoContainer<uint8_t, 256, 16000, 8, uint16_t>;
+    using Hist = cms::cuda::HistoContainer<uint8_t, 256, 16000, 8, uint16_t>;
     __shared__ Hist hist;
     __shared__ typename Hist::Counter hws[32];
     for (auto j = threadIdx.x; j < Hist::totbins(); j += blockDim.x) {
@@ -101,7 +101,7 @@ namespace gpuVertexFinder {
         nn[i]++;
       };
 
-      forEachInBins(hist, izt[i], 1, loop);
+      cms::cuda::forEachInBins(hist, izt[i], 1, loop);
     }
 
     __syncthreads();
@@ -122,7 +122,7 @@ namespace gpuVertexFinder {
         mdist = dist;
         iv[i] = j;  // assign to cluster (better be unique??)
       };
-      forEachInBins(hist, izt[i], 1, loop);
+      cms::cuda::forEachInBins(hist, izt[i], 1, loop);
     }
 
     __syncthreads();
@@ -171,7 +171,7 @@ namespace gpuVertexFinder {
         mdist = dist;
         minJ = j;
       };
-      forEachInBins(hist, izt[i], 1, loop);
+      cms::cuda::forEachInBins(hist, izt[i], 1, loop);
       // should belong to the same cluster...
       assert(iv[i] == iv[minJ]);
       assert(nn[i] <= nn[iv[i]]);

RecoPixelVertexing/PixelVertexFinding/plugins/gpuClusterTracksDBSCAN.h

@@ -44,7 +44,7 @@ namespace gpuVertexFinder {
     assert(pdata);
     assert(zt);
 
-    using Hist = HistoContainer<uint8_t, 256, 16000, 8, uint16_t>;
+    using Hist = cms::cuda::HistoContainer<uint8_t, 256, 16000, 8, uint16_t>;
     __shared__ Hist hist;
     __shared__ typename Hist::Counter hws[32];
     for (auto j = threadIdx.x; j < Hist::totbins(); j += blockDim.x) {
@@ -96,7 +96,7 @@ namespace gpuVertexFinder {
         nn[i]++;
       };
 
-      forEachInBins(hist, izt[i], 1, loop);
+      cms::cuda::forEachInBins(hist, izt[i], 1, loop);
     }
 
     __syncthreads();
@@ -118,7 +118,7 @@ namespace gpuVertexFinder {
         mz = zt[j];
         iv[i] = j;  // assign to cluster (better be unique??)
       };
-      forEachInBins(hist, izt[i], 1, loop);
+      cms::cuda::forEachInBins(hist, izt[i], 1, loop);
     }
 
     __syncthreads();
@@ -172,7 +172,7 @@ namespace gpuVertexFinder {
         }
         assert(iv[i] == iv[j]);
       };
-      forEachInBins(hist, izt[i], 1, loop);
+      cms::cuda::forEachInBins(hist, izt[i], 1, loop);
     }
     __syncthreads();
 #endif
@@ -194,7 +194,7 @@ namespace gpuVertexFinder {
         mdist = dist;
         iv[i] = iv[j];  // assign to cluster (better be unique??)
       };
-      forEachInBins(hist, izt[i], 1, loop);
+      cms::cuda::forEachInBins(hist, izt[i], 1, loop);
     }
 
     __shared__ unsigned int foundClusters;

RecoPixelVertexing/PixelVertexFinding/plugins/gpuClusterTracksIterative.h

@@ -44,7 +44,7 @@ namespace gpuVertexFinder {
     assert(pdata);
     assert(zt);
 
-    using Hist = HistoContainer<uint8_t, 256, 16000, 8, uint16_t>;
+    using Hist = cms::cuda::HistoContainer<uint8_t, 256, 16000, 8, uint16_t>;
     __shared__ Hist hist;
     __shared__ typename Hist::Counter hws[32];
     for (auto j = threadIdx.x; j < Hist::totbins(); j += blockDim.x) {
@@ -97,7 +97,7 @@ namespace gpuVertexFinder {
         nn[i]++;
       };
 
-      forEachInBins(hist, izt[i], 1, loop);
+      cms::cuda::forEachInBins(hist, izt[i], 1, loop);
     }
 
     __shared__ int nloops;
@@ -165,7 +165,7 @@ namespace gpuVertexFinder {
         mdist = dist;
         iv[i] = iv[j];  // assign to cluster (better be unique??)
       };
-      forEachInBins(hist, izt[i], 1, loop);
+      cms::cuda::forEachInBins(hist, izt[i], 1, loop);
     }
 
     __shared__ unsigned int foundClusters;

RecoPixelVertexing/PixelVertexFinding/plugins/gpuVertexFinderImpl.h

@@ -109,7 +109,7 @@ namespace gpuVertexFinder {
     loadTracks<<<numberOfBlocks, blockSize, 0, stream>>>(tksoa, soa, ws_d.get(), ptMin);
     cudaCheck(cudaGetLastError());
 #else
-    cudaCompat::resetGrid();
+    cms::cudacompat::resetGrid();
     init(soa, ws_d.get());
     loadTracks(tksoa, soa, ws_d.get(), ptMin);
 #endif