 }
 
 /*
+ * Histogram smoothing helper
+ */
+template<int D>
+__device__
+inline static float smoothe( const float* const src, const int bin )
+{
+    const int prev = (bin == 0) ? ORI_NBINS-1 : bin-1;
+    const int next = (bin == ORI_NBINS-1) ? 0 : bin+1;
+
+    const float f  = ( src[prev] + src[bin] + src[next] ) / 3.0f;
+
+    return f;
+}
+
+/*
  * Compute the keypoint orientations for each extremum
  * using 16 threads for each of them.
  * direct curve fitting approach

    __syncthreads();

#ifdef WITH_VLFEAT_SMOOTHING
    for( int i=0; i<3; i++ ) {
        for( int bin = threadIdx.x; bin < ORI_NBINS; bin += blockDim.x ) {
            int prev = bin == 0 ? ORI_NBINS-1 : bin-1;
            int next = bin == ORI_NBINS-1 ? 0 : bin+1;
            sm_hist[bin] = ( hist[prev] + hist[bin] + hist[next] ) / 3.0f;
        }
        __syncthreads();
        for( int bin = threadIdx.x; bin < ORI_NBINS; bin += blockDim.x ) {
            int prev = bin == 0 ? ORI_NBINS-1 : bin-1;
            int next = bin == ORI_NBINS-1 ? 0 : bin+1;
            hist[bin] = ( sm_hist[prev] + sm_hist[bin] + sm_hist[next] ) / 3.0f;
        }
        __syncthreads();
    }
    for( int bin = threadIdx.x; bin < ORI_NBINS; bin += blockDim.x ) {
        sm_hist[bin] = hist[bin];
    }
    __syncthreads();
#else // not WITH_VLFEAT_SMOOTHING
    for( int bin = threadIdx.x; bin < ORI_NBINS; bin += blockDim.x ) {

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

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

    hist[threadIdx.x+ 0] = 0.0f;
    hist[threadIdx.x+32] = 0.0f;

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