     }
 };
 
-__shared__ int2  edge_buffer[2][EDGE_LINKING_MAX_EDGE_LENGTH]; // convexEdgeSegment
-__shared__ int   edge_buffer_head[2];
+__shared__ int2  edge_buffer[EDGE_LINKING_MAX_EDGE_LENGTH]; // convexEdgeSegment
+__shared__ int   edge_index[2];
 
 struct EdgeBuffer
 {

        }
        int j = 0;
        int2* ptr = output.ptr(idx);
        for( int i=edge_buffer_head[Right]; i>=0; i-- ) {
            ptr[j] = edge_buffer[Right][i];
            j++;
        }
        for( int i=1; i<=edge_buffer_head[Left]; i++ ) {
            ptr[j] = edge_buffer[Left][i];
            j++;
        }
    }
};


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

    __device__ inline
    void append( Direction d, int2 val )
    {
        if( d == Left ) {
            edge_buffer[edge_index[Left]] = val;
            inc( edge_index[Left] );
            assert( edge_index[Left] != edge_index[Right] );
        } else {
            dec( edge_index[Right] );
            edge_buffer[edge_index[Right]] = val;
            assert( edge_index[Left] != edge_index[Right] );
        }
    }

    __device__ inline
    int size() const
    {
        if( edge_index[Left] > edge_index[Right] ) {
            return edge_index[Left] - edge_index[Right];
        } else {
            return edge_index[Left] + ( EDGE_LINKING_MAX_EDGE_LENGTH - edge_index[Right] );
        }
    }

    __device__
