     ext->num_ori = angles;
 }
 
+class ExtremaRead
+{
+    const Extremum* const _oris;
+public:
+    inline __device__
+    ExtremaRead( const Extremum* const d_oris ) : _oris( d_oris ) { }
+
+    inline __device__
+    int operator()( int n ) const { return _oris[n].num_ori; }
+};
+
+class ExtremaWrt
+{
+    Extremum* _oris;
+public:
+    inline __device__
+    ExtremaWrt( Extremum* d_oris ) : _oris( d_oris ) { }
+
+    inline __device__
+    int& operator()( int n ) const { return _oris[n].idx_ori; }
+};
+
+class ExtremaTot
+{
+    int* _extrema_counter;
+public:
+    inline __device__
+    ExtremaTot( int* extrema_counter ) : _extrema_counter( extrema_counter ) { }
+
+    inline __device__
+    int& operator()() { return *_extrema_counter; }
+};
 
 __global__
-void orientation_starter_v1( Extremum*     extremum,
-                             int*          extrema_counter,
-                             Plane2D_float layer )
-{
-    dim3 block;
-    dim3 grid;
-    grid.x  = *extrema_counter;
-    block.x = ORI_V1_NUM_THREADS;
-
-    if( grid.x != 0 ) {
-        compute_keypoint_orientations_v1
-            <<<grid,block>>>
-            ( extremum,
-              extrema_counter,
-              layer );
+void ori_prefix_sum( int*      extrema_counter,
+                     int*      featvec_counter,
+                     Extremum* extremum )
+{
+    if( threadIdx.x == 0 && threadIdx.y == 0 ) {
+        printf("Enter %s, %d extrema\n", __func__, *extrema_counter );
+    }
+    __syncthreads();
+
+    ExtremaRead r( extremum );
+    ExtremaWrt  w( extremum );
+    ExtremaTot  t( featvec_counter );
+    ExclusivePrefixSum<ExtremaRead,ExtremaWrt,ExtremaTot>( *extrema_counter, r, w, t );
+
+    __syncthreads();
+
+    if( threadIdx.x == 0 && threadIdx.y == 0 ) {
+        *featvec_counter = min( *featvec_counter, d_max_orientations );
+    }
+
+    if( threadIdx.x == 0 && threadIdx.y == 0 ) {
+        printf("Leave %s, %d extrema -> %d oris\n", __func__, *extrema_counter, *featvec_counter );
     }
 }
 

            grid.x  = h_num_extrema[level];
            block.x = ORI_V1_NUM_THREADS;
            if( grid.x != 0 ) {
                    compute_keypoint_orientations_v2
                        <<<grid,block,0,oct_str>>>
                        ( oct_obj.getExtrema( level ),
                          &d_num_extrema[level],
                          oct_obj.getData( level ),
                          &orientation_num_blocks[level],
                          grid.x * grid.y );
            }
        }
    }

 *************************************************************/

__global__
void orientation_starter( Extremum*     extremum,
                             int*          extrema_counter,
                             int*          featvec_counter,
                             Plane2D_float layer )
{
#ifdef USE_DYNAMIC_PARALLELISM // defined in_s_pyramid.h
