Fix race condition

zhaorz · May 11, 2017 · ce56f93 · ce56f93
1 parent 20e3ae5
commit ce56f93
Showing 1 changed file with 9 additions and 0 deletions.
diff --git a/src/kernels/optimize.cu b/src/kernels/optimize.cu
@@ -89,6 +89,8 @@ __device__ void calcProjection(dev_patch_state* states,
 
   }
 
+  __syncthreads();
+
 }
 
 
@@ -117,6 +119,7 @@ __global__ void kernelInterpolateAndComputeErr(
       calcProjection(states, patchX, patchY, tempX, tempY, raw, patch_size,
           out_thresh, lb, ubw, ubh);
     }
+    __syncthreads();
 
 
     // Interpolate the patch
@@ -158,6 +161,8 @@ __global__ void kernelInterpolateAndComputeErr(
 
     }
 
+    __syncthreads();
+
     // Compute mean
     __shared__ float mean;
 
@@ -181,6 +186,8 @@ __global__ void kernelInterpolateAndComputeErr(
       cost[i] = raw[i] * raw[i];
     }
 
+    __syncthreads();
+
     if (tid == 0) {
       float c = 0.0;
       for (int i = 0; i < patch_size * patch_size * 3; i++) {
@@ -216,6 +223,8 @@ __global__ void kernelInterpolateAndComputeErr(
 
     notFirst = true;
 
+    __syncthreads();
+
   }
 
 }