                                  { 0,  1,  1,  1,  0, -1, -1, -1} };
 
 #define EDGE_LINKING_MAX_EDGE_LENGTH    100
-
-template<typename T>
-struct ListReplacement
-{
-    T oops;
-
-    __device__ ListReplacement() { }
-    __device__ void push_back( T& val ) { oops = val; }
-    __device__ void pop_front( ) { }
-    __device__ T&   back()  { return oops; }
-    __device__ T&   front() { return oops; }
-    __device__ int  size() const { return 1; }
+#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 edges         The 0/1 map of edge points
- * @param d_dx
- * @param d_dy
- * @param triplepoints  The array of points including voters and seeds
- * @param seed_indices  The array of indices of seeds in triplepoints
+ * @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( cv::cuda::PtrStepSzb         edges,
-                   DevEdgeList<TriplePoint>     triplepoints,
-                   DevEdgeList<int>             seed_indices,
-                   const cv::cuda::PtrStepSz16s d_dx,
-                   const cv::cuda::PtrStepSz16s d_dy,
-                   size_t param_windowSizeOnInnerEllipticSegment,
-                   float  param_averageVoteMin )
-{
-    // pmax              = one seed
-
-    const int offset = blockIdx.x * 32 + threadIdx.z;
-    int direction = ( threadIdx.y == 0 ) ?  1   // link left
-                                         : -1;  // link right
-
-    if( offset == 0 ) return;
-
-    int idx = seed_indices.ptr[offset];
-    if( idx >= triplepoints.Size() ) return;
-
-    TriplePoint* p = &triplepoints.ptr[idx];
-
-    __shared__ TriplePoint* convexEdgeSegment[2 * EDGE_LINKING_MAX_EDGE_LENGTH+1];
-    __shared__ int          ces_idx[2]; // 0 offset right, 1 negative offset left
-    __shared__ float        averageVoteCollect[2];
-    float                   averageVote =  p->_winnerSize;
+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 ) {
-        ces_idx[threadIdx.y] = EDGE_LINKING_MAX_EDGE_LENGTH + 1 + direction;
-        averageVoteCollect[threadIdx.y] = ( threadIdx.y==0 ) ? averageVote : 0;
-        if( threadIdx.y == 0 ) {
-            convexEdgeSegment[EDGE_LINKING_MAX_EDGE_LENGTH+1] = p;
-        }
+        buf.init( p->coord );
     }
     __syncthreads();
 
-    ListReplacement<float> phi;
-    std::size_t i = 0;
-    bool found = true;
-
-    int stop = 0;
-
+    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[threadIdx.y];
-
-    while( (i < EDGE_LINKING_MAX_EDGE_LENGTH) && (found) && (averageVote >= param_averageVoteMin) )
-    {
-        bool skip;
+    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 ) {
-            skip = atomicExch( &p->edgeLinking.processed, true );
-        }
-        skip = __shfl( skip, 0 );
-        if( skip ) {
-            // Another warp has processed this point or is processing this
-            // point.
-            // End processing.
-            break;
-        }
-
-        // Angle refers to the gradient direction angle (might be optimized):
-        float angle = fmodf( atan2f(p->d.x,p->d.y) + 2.0f * CUDART_PI_F, 2.0f * CUDART_PI_F );
-
-        phi.push_back(angle);
-
-        if (phi.size() > param_windowSizeOnInnerEllipticSegment) // TODO , 4 est un paramètre de l'algorithme, + les motifs à détecter sont importants, + la taille de la fenêtre doit être grande
-        {
-            phi.pop_front();
-        }
+            // 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;

                     ( ( 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;

     * Seeds may have int-valued indices, so we need 4-byte cells.
     * _d_intermediate is a float image, and can be used.
     */
    // cv::cuda::PtrStepSz32s edge_cast( _d_intermediate );

    dim3 block;
    dim3 grid;

