 }
 
 __global__ void descriptor_starter( int*          extrema_counter,
+                                    int*          featvec_counter,
                                     Extremum*     extrema,
                                     Descriptor*   descs,
                                     Plane2D_float layer )


            int* extrema_counters = oct_obj.getExtremaCounterD();
            int* extrema_counter  = &extrema_counters[level];
            descriptor_starter
                <<<1,1,0,oct_str>>>
                ( extrema_counter,
                  oct_obj.getExtrema( level ),
                  oct_obj.getDescriptors( level ),
                  data );

    // it may be good to start more threads, but this kernel
    // is too fast to be noticable in profiling

    grid.x  = grid_divide( *featvec_counter, 32 );
    block.x = 32;
    block.y = 32;
    block.z = 1;

    normalize_histogram
        <<<grid,block>>>
        ( descs, *featvec_counter );
#endif // not USE_DYNAMIC_PARALLELISM
}

