From b6d040bae8b3af16215f1d42f10796a12abd9c4e Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 25 Mar 2020 00:28:04 +0100 Subject: [PATCH] Integrate the comments from the upstream PRs (cms-patatrack#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 --- .../plugins/gpuClusterTracksByDensity.h | 8 ++++---- .../plugins/gpuClusterTracksDBSCAN.h | 10 +++++----- .../plugins/gpuClusterTracksIterative.h | 6 +++--- .../PixelVertexFinding/plugins/gpuVertexFinderImpl.h | 2 +- 4 files changed, 13 insertions(+), 13 deletions(-) diff --git a/RecoPixelVertexing/PixelVertexFinding/plugins/gpuClusterTracksByDensity.h b/RecoPixelVertexing/PixelVertexFinding/plugins/gpuClusterTracksByDensity.h index 871b09599c903..b32c7d5b613db 100644 --- a/RecoPixelVertexing/PixelVertexFinding/plugins/gpuClusterTracksByDensity.h +++ b/RecoPixelVertexing/PixelVertexFinding/plugins/gpuClusterTracksByDensity.h @@ -48,7 +48,7 @@ namespace gpuVertexFinder { assert(pdata); assert(zt); - using Hist = HistoContainer; + using Hist = cms::cuda::HistoContainer; __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]]); diff --git a/RecoPixelVertexing/PixelVertexFinding/plugins/gpuClusterTracksDBSCAN.h b/RecoPixelVertexing/PixelVertexFinding/plugins/gpuClusterTracksDBSCAN.h index 593c7597aecea..ffd7fdc948bf8 100644 --- a/RecoPixelVertexing/PixelVertexFinding/plugins/gpuClusterTracksDBSCAN.h +++ b/RecoPixelVertexing/PixelVertexFinding/plugins/gpuClusterTracksDBSCAN.h @@ -44,7 +44,7 @@ namespace gpuVertexFinder { assert(pdata); assert(zt); - using Hist = HistoContainer; + using Hist = cms::cuda::HistoContainer; __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; diff --git a/RecoPixelVertexing/PixelVertexFinding/plugins/gpuClusterTracksIterative.h b/RecoPixelVertexing/PixelVertexFinding/plugins/gpuClusterTracksIterative.h index 14c825f353960..49da86e941867 100644 --- a/RecoPixelVertexing/PixelVertexFinding/plugins/gpuClusterTracksIterative.h +++ b/RecoPixelVertexing/PixelVertexFinding/plugins/gpuClusterTracksIterative.h @@ -44,7 +44,7 @@ namespace gpuVertexFinder { assert(pdata); assert(zt); - using Hist = HistoContainer; + using Hist = cms::cuda::HistoContainer; __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; diff --git a/RecoPixelVertexing/PixelVertexFinding/plugins/gpuVertexFinderImpl.h b/RecoPixelVertexing/PixelVertexFinding/plugins/gpuVertexFinderImpl.h index 10c487d588e9d..0da24cef219e0 100644 --- a/RecoPixelVertexing/PixelVertexFinding/plugins/gpuVertexFinderImpl.h +++ b/RecoPixelVertexing/PixelVertexFinding/plugins/gpuVertexFinderImpl.h @@ -109,7 +109,7 @@ namespace gpuVertexFinder { loadTracks<<>>(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