 #include "assist.h"
 #include "recursive_sweep.h"
 
-#define EDGE_NOT_FOUND -1
-#define CONVEXITY_LOST -2
-#define LOW_FLOW -3
-
 using namespace std;
 
 namespace popart
 {
 
+namespace linking
+{
 __constant__
 static int xoff_select[8]    =   { 1,  1,  0, -1, -1, -1,  0,  1};
 

                                 { 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; }
};

/**
 * @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 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;
    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;
        }
    }
    __syncthreads();

    ListReplacement<float> phi;
    std::size_t i = 0;
    bool found = true;

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

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

     * This label is their index in the _seed_indices list, because
     * it is a unique int strictly > 0
     */
    block.x = 8;
    block.y = 0;
    block.z = 1;
    grid.x  = _vote._seed_indices.host.size;
    grid.y  = 1;
    grid.z  = 1;

    linking::edge_linking
        <<<grid,block,0,_stream>>>
        ( _vote._chained_edgecoords.dev,
          _vote._seed_indices.dev,
          _vote._d_edgepoint_index_table,
          _d_edges,
          _d_ring_counter,
          _d_ring_output,
          param.windowSizeOnInnerEllipticSegment,
          param.averageVoteMin );

