         }
         int j = 0;
         int2* ptr = output.ptr(idx);
-        for( int i=edge_buffer_head[Right]; i>=0; i-- ) {
-            ptr[j] = edge_buffer[Right][i];
+        for( int idx=edge_index[Right]; idx!=edge_index[Left]; inc(idx) ) {
+            ptr[j] = edge_buffer[idx];
             j++;
         }
-        for( int i=1; i<=edge_buffer_head[Left]; i++ ) {
-            ptr[j] = edge_buffer[Left][i];
-            j++;
-        }
+    }
+
+private:
+    __device__ inline void inc( size_t& idx )
+    {
+        idx = ( idx >= EDGE_LINKING_MAX_EDGE_LENGTH-1 ) ? 0 : idx + 1;
+    }
+
+    __device__ inline void dec( size_t& idx )
+    {
+        idx = ( idx == 0 ) ? EDGE_LINKING_MAX_EDGE_LENGTH-1 : idx - 1;
     }
 };
 

        // This direction still has points.
        // We can identify the highest threadId / lowest rotation value j
        uint32_t computer = __ffs( any_point_found ) - 1;
        if( computer >= 8 ) {
            printf("point found by thread %d\n", computer);
            assert( computer < 8 );
        }

        found = LOW_FLOW;


    void init( int2 start )
    {
        if( threadIdx.x == 0 ) {
            edge_buffer[0]    = start;
            edge_index[Left]  = 1;
            edge_index[Right] = 0;
        }
        __syncthreads();
    }
