Work around atomicAdd synchronisation problem on Volta (#180)

cms-patatrack · Nov 16, 2020 · be8f851 · be8f851
1 parent 2092143
commit be8f851
Showing 1 changed file with 2 additions and 2 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);
         }
       }