             //
             float s;
             float c;
-            __sincosf( phi.back() - phi.front(), &s, &c );
+            __sincosf( phi.diff(), &s, &c );
             s = ( direction == Left ) ? s : -s;
 
             //

            }
        }

        float angle;

        if( threadIdx.x == 0 ) {
            // this cycle gradients
            short2 tcg;

            // Angle refers to the gradient direction angle (might be optimized):
            tcg.x = d_dx.ptr(tcc.y)[tcc.y];
            tcg.y = d_dy.ptr(tcc.y)[tcc.y];
            float atanval = atan2f( tcg.x, tcg.y );

            angle = fmodf( atanval + 2.0f * CUDART_PI_F, 2.0f * CUDART_PI_F );

            phi.push_back( angle );
        }
        angle = __shfl( angle, 0 );

        int shifting = rintf( ( (angle + CUDART_PI_F / 4.0f)
                              / (2.0f * CUDART_PI_F) ) * 8.0f ) - 1;

    FULL_CIRCLE         =  3
};

struct RingBuffer
{
    Direction direction;
    float     ring_buffer[EDGE_LINKING_MAX_RING_BUFFER_SIZE];
    size_t    max_size;
    size_t    front_index;
    size_t    back_index;
    size_t    ct;
    float     angle_diff;

    __device__ RingBuffer( Direction dir, size_t size ) {
        assert( size < EDGE_LINKING_MAX_RING_BUFFER_SIZE );
        direction   = dir;
        max_size    = size;
        front_index = 0;
        back_index  = 0;
        ct          = 0;
        angle_diff  = 0.0f;
    }

    __device__ inline void set_direction( Direction dir ) {
        direction = dir;
    }

    __device__ void push_back( float angle )
    {
        if( threadIdx.x == 0 ) {
            if( direction == Left ) {
                ring_buffer[back_index] = angle;
                inc( back_index );
                if( front_index == back_index )
                    inc( front_index );
                else
                    ct++:
            } else {
                dec( front_index );
                if( front_index == back_index )
                    dec( back_index );
                else
                    ct++;
                ring_buffer[front_index] = angle;
            }
            angle_diff = back() - front();
        }
        angle_diff = __shfl( angle_diff, 0 );
        ct         = __shfl( ct, 0 );
    }

    __device__ inline float diff( )
    {
        // replace in lilian's original code back()-front()
        return angle_diff;
    }

    __device__ int  inline size() const {
        return ct;
    }
private:
    __device__ inline float front( )
    {
        assert( threadIdx.x == 0 );
        assert( front_index != back_index );
        return ring_buffer[direction][front_index];
    }

    __device__ inline float back( )
    {
        assert( threadIdx.x == 0 );
        assert( front_index != back_index );
        size_t lookup = back_index;
        dec( lookup );
        return ring_buffer[direction][lookup];
    }

    __device__ inline void inc( size_t& idx )
    {
        assert( threadIdx.x == 0 );
        idx = ( idx >= max_size-1 ) ? 0 : idx + 1;
    }

    __device__ inline void dec( size_t& idx )
    {
        assert( threadIdx.x == 0 );
        idx = ( idx == 0 ) ? max_size-1 : idx - 1;
    }
};

