     cerr << "Enter " << __FUNCTION__ << endl;
 
     _keep_time_orient_v1.start();
+compute_keypoint_orientations_v1
+                    <<<dim3(16),dim3(1)>>>
+                    ( _octaves[octave].getExtrema( level ),
+                      _octaves[octave].getExtremaMgmtD( ),
+                      level,
+                      _octaves[octave].getData( level ) );
+/*
     for( int octave=0; octave<_num_octaves; octave++ ) {
         _octaves[octave].readExtremaCount( );
         cudaDeviceSynchronize( );

            }
        }
    }
    _keep_time_orient_v1.stop();

    cerr << "Leave " << __FUNCTION__ << endl;

        hist[threadIdx.x  ] = 0.25f * prev1 + 0.5f * tmp1_curr + 0.25f * tmp1_next;
        hist[threadIdx.x*2] = 0.25f * prev2 + 0.5f * tmp2_curr + 0.25f * tmp2_next;

        //printf("(%d): bin %d: %f\n", threadIdx.x, hist[threadIdx.x]);
        //printf("(%d): bin %d: %f\n", threadIdx.x*2, hist[threadIdx.x*2]);

        //0->31 done
        //missing 32->35 (4 bins)
        if(threadIdx.x < 3){
	//0+32
	//1+32
	//2+32 34
            //bin 32->34
            float prev = hist[(threadIdx.x+32 - 1 + ORI_V1_NUM_THREADS) % ORI_V1_NUM_THREADS];
            float curr = hist[threadIdx.x]+32; //32->34
            float next = hist[threadIdx.x+33]; //33->35
            hist[threadIdx.x +32] = 0.25f * prev + 0.5f * curr + 0.25f * next;
            //printf("(%d): bin %d: %f\n", threadIdx.x+32, hist[threadIdx.x+32]);
        }else if(threadIdx.x == 3){
            //bin 35
            hist[threadIdx.x +32] = 0.25f * hist[threadIdx.x +32 -1]
                                   + 0.5f * hist[threadIdx.x +32] + 0.25f * first;
            //printf("(%d): bin %d: %f\n", threadIdx.x+32, hist[threadIdx.x+32]);
        }
    }
#else
