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