     inline __device__
     bool first_contrast_ok( const float val ) const;
 
-    inline __device__
-    bool refine( const float3& d, int3& n, const int width, const int height, const int maxlevel );
+    /* refine
+     * returns -1 : break loop and fail
+     *          0 : continue looping
+     *          1 : break loop and succeed
+     */
+    inline __device__
+    int refine( float3& d, int3& n, const int width, const int height, const int maxlevel, bool last_it );
+
+    /*
+     * returns true  : values after refine make sense
+     *         false : they do not
+     */
+    inline __device__
+    bool verify( const float xn, const float yn, const float sn, const int width, const int height, const int maxlevel ) const;
+};
+
+template<>
+class ModeFunctions<Config::OpenCV>
+{
+public:
+    inline __device__
+    bool first_contrast_ok( const float val ) const
+    {
+        return ( fabsf( val ) >= floorf( d_threshold ) );
+    }
+
+    inline __device__
+    int refine( float3& d, int3& n, const int width, const int height, const int maxlevel, bool last_it ) const
+    {
+        // OpenCV mode is a special case because d remain unmodified.
+        // Either we return 1, and n has not been modified.
+        // Or we quit the loop by exceeding the limit, and reject the point anyway.
+
+        float3 t = make_float3( fabsf(d.x), fabsf(d.y), fabsf(d.z) );
+
+        if( t.x < 0.5f && t.y < 0.5f && t.z < 0.5f ) {
+            // return false, quit the loop, success
+            return 1;
+        }
+
+        // This test in OpenCV is totally useless in CUDA because the thread
+        // would simple idle before failing 7 instructions below anyway.
+        // if( t.x > (float)(INT_MAX/3) || t.y > (float)(INT_MAX/3) || t.z > (float)(INT_MAX/3) ) {
+            // return false, quit the loop, fail
+            // return -1;
+        // }
+
+        n.x += rintf( d.x );  // quicker roundf
+        n.y += rintf( d.y );  // quicker roundf
+        n.z += rintf( d.z );  // quicker roundf
+
+        const int retval = ( n.x < 5 || n.x >= width-5 ||
+                             n.y < 5 || n.y >= height-5 ||
+                             n.z < 1 || n.z > maxlevel-2 ) ? -1 : 0;
+            // if outside of all DoG images (minus border),
+            // quit the loop, fail
+
+        return retval;
+    }
+
+    inline __device__
+    int verify( const float xn, const float yn, const float sn, const int width, const int height, const int maxlevel ) const
+    {
+        return true;
+        // return ( xn < 5 || xn >= width-5 || yn < 5 || yn >= height-5 || sn < 1 || sn > maxlevel-2 ) ? false : true;
+    }
 };
 
 template<>

     */
    const int block_x = blockIdx.x * 32;
    const int block_y = blockIdx.y * blockDim.y;
    const int y       = block_y + threadIdx.y;
    const int x       = block_x + threadIdx.x;

    // int32_t x0 = x;
    // int32_t x1 = x+1;

    /* negative determinant => curvatures have different signs -> reject it */
    if (det <= 0.0f) {
#ifdef PRINT_EXTREMA_DEBUG_INFO
        printf("Found an extremum at %d %d (o=%d,l=%d) - negative determinant\n", x, y, debug_octave, level );
#endif // PRINT_EXTREMA_DEBUG_INFO
        return false;
    }

    /* accept-reject extremum */
    // if( fabsf(contr) < (d_threshold*2.0f) )
    if( fabsf(contr) < scalbnf( d_threshold, 1 ) )
    {
#ifdef PRINT_EXTREMA_DEBUG_INFO
        printf("Found an extremum at %d %d (o=%d,l=%d) - 2nd peak tresh failed\n", x, y, debug_octave, level );
#endif // PRINT_EXTREMA_DEBUG_INFO
        return false;
    }
