Work around atomicAdd synchronisation problem on Volta (#180)

cms-patatrack · Sep 30, 2018 · e269bdd · e269bdd
1 parent b5135e6
commit e269bdd
Show file tree

Hide file tree

Showing 2 changed files with 4 additions and 4 deletions.
diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h
@@ -198,7 +198,7 @@ namespace gpuClustering {
         printf("# loops %d\n",nloops);
 #endif
 
-    __shared__ int foundClusters;
+    __shared__ unsigned int foundClusters;
     foundClusters = 0;
     __syncthreads();
 
@@ -208,7 +208,7 @@ namespace gpuClustering {
         if (id[i] == InvId)                 // skip invalid pixels
           continue;
         if (clusterId[i] == i) {
-          auto old = atomicAdd(&foundClusters, 1);
+          auto old = atomicInc(&foundClusters, 0xffffffff);
           clusterId[i] = -(old + 1);
         }
       }

diff --git a/RecoPixelVertexing/PixelVertexFinding/src/gpuClusterTracks.h b/RecoPixelVertexing/PixelVertexFinding/src/gpuClusterTracks.h
@@ -184,7 +184,7 @@ namespace gpuVertexFinder {
     }
 
 
-    __shared__ int foundClusters;
+    __shared__ unsigned int foundClusters;
     foundClusters = 0;
     __syncthreads();
 
@@ -193,7 +193,7 @@ namespace gpuVertexFinder {
     for (int i = threadIdx.x; i < nt; i += blockDim.x) {
       if (iv[i] == i) {
 	if  (nn[i]>=minT) {
-	  auto old = atomicAdd(&foundClusters, 1);
+	  auto old = atomicInc(&foundClusters, 0xffffffff);
 	  iv[i] = -(old + 1);
 	  zv[old]=0;
 	  wv[old]=0;