 // Excellent easy speedup possible by using threadIdx.x for parallel checking
 // of the voter array in count_winners!!!
 
+#define CONC_POINTS 8
+
 namespace popart {
 
 namespace vote {
 
 __device__ inline
-int count_winners( FrameMetaPtr&                   meta,
-                   const int                       inner_point_index,
-                   TriplePoint*                    inner_point,
-                   const DevEdgeList<TriplePoint>& voters,
-                   const float*                    chosen_flow_length,
-                   const DevEdgeList<int>          chosen_idx )
-{
-    int   winner_size = 0;
-    float flow_length = 0.0f;
+void count_winners( FrameMetaPtr&                   meta,
+                    const int                       offset,
+                    const int                       maxpoints,
+                    const int*                      inner_points,
+                    const DevEdgeList<TriplePoint>& voters,
+                    const float*                    chosen_flow_length,
+                    const DevEdgeList<int>          chosen_idx )
+{
+    int          inner_point_index[CONC_POINTS];
+    TriplePoint* inner_point[CONC_POINTS];
+    int          winner_size[CONC_POINTS];
+    float        flow_length[CONC_POINTS];
+
+    for( int point=0; point<CONC_POINTS; point++ ) {
+        const bool pred          = ( offset+point < maxpoints );
+        const int  ipi           = pred ? inner_points[offset+point] : 0;
+        inner_point_index[point] = ipi;
+        inner_point[point]       = pred ? &voters.ptr[ipi] : 0;
+        winner_size[point]       = 0;
+        flow_length[point]       = 0.0f;
+    }
 
     /* This loop looks dangerous, but it is actually faster than
      * a manually partially unrolled loop.
      */
     const int voter_list_size = meta.list_size_voters();
-    for( int i=0; i<voter_list_size; i++ ) {
-        int my_vote = chosen_idx.ptr[i];
-        if( my_vote == inner_point_index ) {
-            winner_size += 1;
-            // flow_length += voters.ptr[i].chosen_flow_length;
-            flow_length += chosen_flow_length[i];
+    const int start_offset = 32 * threadIdx.y + threadIdx.x;
+    const int inc_offset   = 32*32;
+    for( int i=start_offset; i<voter_list_size; i+=inc_offset ) {
+        const int my_vote = chosen_idx.ptr[i];
+        for( int point=0; point<CONC_POINTS; point++ ) {
+            if( inner_point[point] == 0 ) continue;
+
+            if( my_vote == inner_point_index[point] ) {
+                winner_size[point] += 1;
+                flow_length[point] += chosen_flow_length[i];
+            }
         }
     }
-    inner_point->_winnerSize = winner_size;
-    inner_point->_flowLength = flow_length / winner_size;
-    return winner_size;
+
+    for( int point=0; point<CONC_POINTS; point++ ) {
+        if( inner_point[point] == 0 ) continue;
+
+        winner_size[point] += __shfl_down( winner_size[point], 16 );
+        winner_size[point] += __shfl_down( winner_size[point],  8 );
+        winner_size[point] += __shfl_down( winner_size[point],  4 );
+        winner_size[point] += __shfl_down( winner_size[point],  2 );
+        winner_size[point] += __shfl_down( winner_size[point],  1 );
+        winner_size[point]  = __shfl     ( winner_size[point],  0 );
+
+        flow_length[point] += __shfl_down( flow_length[point], 16 );
+        flow_length[point] += __shfl_down( flow_length[point],  8 );
+        flow_length[point] += __shfl_down( flow_length[point],  4 );
+        flow_length[point] += __shfl_down( flow_length[point],  2 );
+        flow_length[point] += __shfl_down( flow_length[point],  1 );
+        flow_length[point]  = __shfl     ( flow_length[point],  0 );
+    }
+
+    __shared__ int   winner_array[CONC_POINTS][32];
+    __shared__ float length_array[CONC_POINTS][32];
+
+    if( threadIdx.x < CONC_POINTS ) {
+        const int point = threadIdx.x;
+
+        winner_array[point][threadIdx.y] = winner_size[point];
+        length_array[point][threadIdx.y] = flow_length[point];
+    }
+
+    __syncthreads();
+
+    if( threadIdx.y < CONC_POINTS ) {
+        const int point = threadIdx.y;
+
+        if( inner_point[point] == 0 ) return;
+
+        winner_size[point] = winner_array[point][threadIdx.x];
+        flow_length[point] = length_array[point][threadIdx.x];
+
+        winner_size[point] += __shfl_down( winner_size[point], 16 );
+        winner_size[point] += __shfl_down( winner_size[point],  8 );
+        winner_size[point] += __shfl_down( winner_size[point],  4 );
+        winner_size[point] += __shfl_down( winner_size[point],  2 );
+        winner_size[point] += __shfl_down( winner_size[point],  1 );
+
+        flow_length[point] += __shfl_down( flow_length[point], 16 );
+        flow_length[point] += __shfl_down( flow_length[point],  8 );
+        flow_length[point] += __shfl_down( flow_length[point],  4 );
+        flow_length[point] += __shfl_down( flow_length[point],  2 );
+        flow_length[point] += __shfl_down( flow_length[point],  1 );
+
+        if( threadIdx.x == 0 ) {
+            inner_point[point]->_winnerSize = winner_size[point];
+            inner_point[point]->_flowLength = flow_length[point] / winner_size[point];
+        }
+    }
 }
 
 /* For all chosen inner points, compute the average flow length and the


__global__
void dp_call_eval_chosen( FrameMetaPtr             meta,
                          DevEdgeList<int>         inner_points,      // input
                          DevEdgeList<TriplePoint> voters, // modified input
                          const float*             chosen_flow_length, // input
                          DevEdgeList<int>         chosen_idx ) // input
{
    int listsize = meta.list_size_inner_points();

    dim3 block( 32, 1, 1 );
    dim3 grid( grid_divide( listsize, 32 ), 1, 1 );

    vote::eval_chosen
        <<<grid,block>>>
        ( meta,
          inner_points,
          voters,
          chosen_flow_length,
          chosen_idx );
}

} // namespace vote

     * value d_num_selected_out from the device before the voting
     * step.
     */
// #ifndef NDEBUG
#if 0
    _voters.copySizeFromDevice( _stream, EdgeListCont );
#endif
    _interm_inner_points.copySizeFromDevice( _stream, EdgeListWait );

// #ifndef NDEBUG
#if 0
    cerr << "Debug voting (without separable compilation)"
         << " # inner points: " << _interm_inner_points.host.size
         << " # voters      : " << _voters.host.size << endl;
