     int hy = ymax - ymin + 1;
     int loops = wx * hy;
 
+    __syncthreads();
     for( int i = threadIdx.x; popsift::any(i < loops); i += blockDim.x )
     {
         if( i < loops ) {

    const int              iext_off =  dobuf.i_ext_off[octave][extremum_index];
    const InitialExtremum* iext     = &dobuf.i_ext_dat[octave][iext_off];

    __shared__ float hist   [ORI_NBINS];
    __shared__ float sm_hist[ORI_NBINS];
    __shared__ float refined_angle[64];
    __shared__ float yval         [64];

    for( int i = threadIdx.x; i < ORI_NBINS; i += blockDim.x )  hist[i] = 0.0f;
    __syncthreads();

    /* keypoint fractional geometry */
    const float x     = iext->xpos;

    __syncthreads();

#ifdef WITH_VLFEAT_SMOOTHING
    for( int i=0; i<3 ; i++ )
    {
        sm_hist[threadIdx.x+ 0] = smoothe<0>( hist, threadIdx.x+ 0 );
        sm_hist[threadIdx.x+32] = smoothe<1>( hist, threadIdx.x+32 );
        __syncthreads();
        hist[threadIdx.x+ 0]    = smoothe<2>( sm_hist, threadIdx.x+ 0 );
        hist[threadIdx.x+32]    = smoothe<3>( sm_hist, threadIdx.x+32 );
        __syncthreads();
    }

    sm_hist[threadIdx.x+ 0] = hist[threadIdx.x+ 0];
    sm_hist[threadIdx.x+32] = hist[threadIdx.x+32];
    __syncthreads();
#else // not WITH_VLFEAT_SMOOTHING
    for( int bin = threadIdx.x; bin < ORI_NBINS; bin += blockDim.x ) {
