      * This label is their index in the _seed_indices list, because
      * it is a unique int strictly > 0
      */
-    block.x = 32;
-    block.y = 1;
+    block.x = 8;
+    block.y = 0;
     block.z = 1;
-    grid.x  = grid_divide( _vote._seed_indices.host.size, 32 );
+    grid.x  = _vote._seed_indices.host.size;
     grid.y  = 1;
     grid.z  = 1;
 
-    edge_linking
+    linking::edge_linking
         <<<grid,block,0,_stream>>>
-        ( _d_edges,
-          _vote._chained_edgecoords.dev,
+        ( _vote._chained_edgecoords.dev,
           _vote._seed_indices.dev,
-          _d_dx,
-          _d_dy,
+          _vote._d_edgepoint_index_table,
+          _d_edges,
+          _d_ring_counter,
+          _d_ring_output,
           param.windowSizeOnInnerEllipticSegment,
           param.averageVoteMin );
 

                     ( ( s <  0.0f   ) && ( c < 0.0f ) ) );
            
            if( not stop ) {
                // this outcome of this test does not depend on
                // threadIdx.x / j

                p->edgeLinking.processed = true;

                convexEdgeSegment[ces_idx[threadIdx.y]] = f;
                ces_idx[threadIdx.y]     += direction;
                averageVoteCollect[threadIdx.y] += f->_winnerSize;
                stop = 1;

                if( f->edgeLinking.processed ) {
                    found = false;
                }
            } else {
                stop  = CONVEXITY_LOST;
                found = false;
            }
        } // end of asynchronous block

        p     = (TriplePoint*)__shfl( (size_t)f,     computer );
        found = __shfl( found, computer );
        ++i;

        __syncthreads();

        averageVote = averageVoteCollect[0] + averageVoteCollect[1];
        int convexEdgeSegmentSize = ces_idx[0] - ces_idx[1];
        averageVote /= convexEdgeSegmentSize;
    } // while


    if( threadIdx.x == 0 && threadIdx.y == 0 )
    {
        int n = 0;
        if ((i == EDGE_LINKING_MAX_EDGE_LENGTH) || (stop == CONVEXITY_LOST)) {
            int convexEdgeSegmentSize = ces_idx[0] - ces_idx[1];
            if (convexEdgeSegmentSize > param_windowSizeOnInnerEllipticSegment) {
                for( int i=ces_idx[1]; i<ces_idx[0]; i++ ) {
                    TriplePoint* collectedP = convexEdgeSegment[i];
                    if (n == convexEdgeSegmentSize - param_windowSizeOnInnerEllipticSegment) {
                        break;
                    } else {
                        // collectedP->_processedIn = true;
                        ++n;
                    }
                }
            }
        } else if (stop == EDGE_NOT_FOUND) {
            for( int i=ces_idx[1]; i<ces_idx[0]; i++ ) {
                TriplePoint* collectedP = convexEdgeSegment[i];
                // collectedP->_processedIn = true;
            }
        }
    }
}

__host__
void Frame::applyLink( const cctag::Parameters& params )
{
    cout << "Enter " << __FUNCTION__ << endl;

    if( _vote._seed_indices.host.size <= 0 ) {
        cout << "Leave " << __FUNCTION__ << endl;

                                 { 0,  1,  1,  1,  0, -1, -1, -1} };

#define EDGE_LINKING_MAX_EDGE_LENGTH    100
#define MAX_RING_BUFFER_SIZE            40

enum Direction
{
    Left = 0,
    Right = 1
};

enum StopCondition
{
    LOW_FLOW        =     -3,
    CONVEXITY_LOST      = -2,
    EDGE_NOT_FOUND      = -1,
    NONE                =  0,
    STILL_SEARCHING     =  1,
    FOUND_NEXT          =  2,
    FULL_CIRCLE         =  3
};

__shared__ float ring_buffer[2][MAX_RING_BUFFER_SIZE];

struct RingBuffer
{
    Direction direction;
    size_t    max_size;
    size_t    front_index;
    size_t    back_index;
    size_t    ct;

    __device__ RingBuffer( Direction dir, size_t size ) {
        direction   = dir;
        max_size    = size;
        front_index = 0;
        back_index  = 0;
        ct          = 0;
    }

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

    __device__ inline size_t inc( size_t& idx )
    {
        return ( idx == max_size-1 ) ? 0 : idx + 1;
    }

    __device__ inline size_t dec( size_t& idx )
    {
        return ( idx == 0 ) ? max_size-1 : idx - 1;
    }

    __device__ void push_back( float angle )
    {
        ring_buffer[direction][front_index].angle = angle;
        back_index = inc( back_index );
        if( front_index == back_index )
            front_index = inc( front_index );
        else
            ct++;
    }

    __device__ inline float front( )
    {
        assert( front_index != back_index );
        return ring_buffer[direction][front_index].angle;
    }

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

    __device__ int  inline size() const {
        return ct;
    }
};

__shared__ int2  edge_buffer[2][EDGE_LINKING_MAX_EDGE_LENGTH]; // convexEdgeSegment

struct EdgeBuffer
{
    int head[2];

    __device__ inline
    void init( int2 start )
    {
        edge_buffer[Left ][0] = start;
        edge_buffer[Right][0] = start;
        head[Left ] = 0;
        head[Right] = 0;
    }

    __device__ inline
    const int2& get( Direction d )
    {
        return edge_buffer[d][head[d]];
    }

    __device__ inline
    void append( Direction d, int2 val )
    {
        head[d]++;
        edge_buffer[d][head[d]] = val;
    }

    __device__ inline
    int size() const
    {
        return head[Left] + head[Right] + 1;
    }

    __device__
    void copy( cv::cuda::PtrStepSzInt2 output, int idx )
    {
        assert( idx < output.rows );
        assert( size() < output.cols );
        int j = 0;
        int2* ptr = output.ptr(idx);
        for( int i=head[Right]; i>=0; i-- ) {
            ptr[j] = edge_buffer[Right][i];
            j++;
        }
        for( int i=1; i<=head[Left]; i++ ) {
            ptr[j] = edge_buffer[Left][i];
            j++;
        }
    }
};

/**
 * @param p Starting seed
 * @param triplepoints          List of all edge points that were potential voters
 * @param edgepoint_index_table Map: map coordinate to triplepoints entry
 * @param edges                 Map: non-0 if edge
 * @param param_windowSizeOnInnerEllipticSegment
 * @param param_averageVoteMin
 */
__device__
void edge_linking_seed( const TriplePoint*           p,
                        DevEdgeList<TriplePoint>     triplepoints,
                        cv::cuda::PtrStepSz32s       edgepoint_index_table, // coord->triplepoint
                        cv::cuda::PtrStepSzb         edges,
                        int*                         d_ring_counter,
                        cv::cuda::PtrStepSzInt2      d_ring_output,
                        size_t param_windowSizeOnInnerEllipticSegment,
                        float  param_averageVoteMin )
{
    const Direction direction       = Left;
    const Direction other_direction = Right;

    float averageVote        = p->_winnerSize;
    float averageVoteCollect = averageVote;

    EdgeBuffer buf;

    if( threadIdx.x == 0 ) {
        buf.init( p->coord );
    }
    __syncthreads();

    RingBuffer<float> phi( Left, param_windowSizeOnInnerEllipticSegment );
    size_t        i     = 0;
    StopCondition found = STILL_SEARCHING;
    // int           stop  = 0;

    const int* xoff = xoff_select;
    const int* yoff = yoff_select[direction];

    // int2   this_cycle_coord = p->coord;
    // short2 this_cycle_gradients;

    while( (i < EDGE_LINKING_MAX_EDGE_LENGTH) && (found==STILL_SEARCHING) && (averageVote >= param_averageVoteMin) )
    {
        // this cycle coordinates
        const int2&  tcc = buf.get( direction );

        if( i > 0 ) {
            // other direction coordinates
            const int2& odc = buf.get( other_direction );

            if( odc.x == tcc.x && odc.y == tcc.y ) {
                // We have gone a full circle.
                // End processing.
                found = FULL_CIRCLE;
                continue;
            }
        }

        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;
