 
             float newbin = __fdividef( num, denB ); // * M_PI/18.0f; // * 10.0f;
             if( newbin >= 0 && newbin <= 2 ) {
-                xcoord[bin] = prev + newbin;
+                refined_angle[bin] = prev + newbin;
                 yval[bin]   = -(num*num) / (4.0f * denB) + hist[prev];
 
 #ifdef LOWE_ORIENTATION_MAX

    __syncthreads();

    // sub-cell refinement of the histogram cell index, yielding the angle
    __shared__ float xcoord[ORI_NBINS];
    __shared__ float yval[ORI_NBINS];

    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;

        bool predicate = ( sm_hist[bin] > max( sm_hist[prev], sm_hist[next] ) );

        const float num  = predicate ? 3.0f *   sm_hist[prev] - 4.0f * sm_hist[bin] + sm_hist[next]   : 0.0f;
        const float denB = predicate ? 2.0f * ( sm_hist[prev] - 2.0f * sm_hist[bin] + sm_hist[next] ) : 1.0f;

        const float newbin = __fdividef( num, denB );

        predicate   = ( predicate && newbin >= 0.0f && newbin <= 2.0f );

        xcoord[bin] = predicate ? prev + newbin : bin;
        yval[bin]   = predicate ?  -(num*num) / (4.0f * denB) + sm_hist[prev] : -INFINITY;
    }
    __syncthreads();

    int best_index = threadIdx.x;
    for( int bin=threadIdx.x+blockDim.x; bin < ORI_NBINS; bin += blockDim.x ) {
        const bool pred = ( yval[bin] > yval[best_index] );
        best_index = pred ? bin : best_index;
    }
    __syncthreads();

    // bitonic sort
    // only for a warp with 32 elements
    for( int i=0; i<5; i++ ) {
        for( int j=i; j>=0; j-- ) {
            const bool swap_dir     = ( threadIdx.x & (2<<i) );
            const int  xorval         = 1 << j;
            const int  other_index  = __shfl_xor( best_index, xorval );
            const bool id_below     = ( threadIdx.x < ( threadIdx.x ^ xorval ) );
            const bool other_bigger = swap_dir ^ ( yval[other_index] > yval[best_index] );
            const bool do_swap      = ( id_below ^ other_bigger == 0 );
            best_index              = do_swap ? other_index : best_index;
        }
    }

    // All threads retrieve the yval of thread 0, the largest
    // of all yvals.
    const float best_val = yval[best_index];
    const float yval_ref = 0.8f * __shfl( best_val, 0 );
    const bool  valid    = ( yval[best_index]  >= yval_ref );
    bool        written  = false;

    if( threadIdx.x < ORIENTATION_MAX_COUNT ) {
        if( valid ) {
            float chosen_bin = xcoord[best_index];
            if( chosen_bin >= ORI_NBINS ) chosen_bin -= ORI_NBINS;
            float th = __fdividef(M_PI2 * chosen_bin , ORI_NBINS) - M_PI;
            ext->orientation[threadIdx.x] = th;


    if(threadIdx.x != 0) return;

    float refined_angle[ORI_NBINS];
    float yval[ORI_NBINS];

    int   maxbin[ORIENTATION_MAX_COUNT];
