Merge pull request #82 from andrew-hardin/fix/two-races

Resolve race conditions identified by racecheck.
alicevision · Apr 20, 2020 · bf3556f · bf3556f
2 parents 42bda27 + f603edf
commit bf3556f
Show file tree

Hide file tree

Showing 2 changed files with 4 additions and 6 deletions.
diff --git a/src/popsift/common/excl_blk_prefix_sum.h b/src/popsift/common/excl_blk_prefix_sum.h
@@ -132,19 +132,15 @@ class Block
                  */
                 _mapping_writer.set( ebs, self, cell );
             }
+            __syncthreads();
 
             if( threadIdx.y == 0 && threadIdx.x == 31 ) {
                 loop_total += ibs;
             }
             __syncthreads();
         }
 
-        // if( threadIdx.y == 0 && threadIdx.x == 31 )
-        if( threadIdx.y == 0 )
-        {
-            loop_total = popsift::shuffle( loop_total, 31 );
-            _total_writer.set( loop_total );
-        }
+        _total_writer.set( loop_total );
     }
 };
 

diff --git a/src/popsift/s_orientation.cu b/src/popsift/s_orientation.cu
@@ -76,6 +76,7 @@ void ori_par( const int           octave,
     __shared__ float sm_hist[ORI_NBINS];
 
     for( int i = threadIdx.x; i < ORI_NBINS; i += blockDim.x )  hist[i] = 0.0f;
+    __syncthreads();
 
     /* keypoint fractional geometry */
     const float x     = iext->xpos;
@@ -204,6 +205,7 @@ void ori_par( const int           octave,
         refined_angle[bin] = predicate ? prev + newbin : -1;
         yval[bin]          = predicate ?  -(num*num) / (4.0f * denB) + sm_hist[prev] : -INFINITY;
     }
+    __syncthreads();
 
     int2 best_index = make_int2( threadIdx.x, threadIdx.x + 32 );