     return true;
 }
 
+template<int sift_mode>
+class ModeFunctions
+{
+public:
+    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 );
+};
+
+template<>
+class ModeFunctions<Config::VLFeat>
+{
+public:
+    inline __device__
+    bool first_contrast_ok( const float val ) const
+    {
+        return ( fabs( val ) >= 0.8 * 2.0 * d_threshold );
+    }
+
+    inline __device__
+    bool refine( const float3& d, int3& n, const int width, const int height, const int maxlevel )
+    {
+        float3 t;
+
+        t.x = ((d.x >= 0.6f && n.x < width-2) ?  1 : 0 )
+            + ((d.x <= -0.6f && n.x > 1)? -1 : 0 );
+
+        t.y = ((d.y >= 0.6f && n.y < height-2)  ?  1 : 0 )
+            + ((d.y <= -0.6f && n.y > 1) ? -1 : 0 );
+
+        // t.z = ((d.z >= 0.6f && n.z < maxlevel-1)  ?  1 : 0 )
+            // + ((d.z <= -0.6f && n.z > 1) ? -1 : 0 );
+
+        if( t.x == 0 && t.y == 0 ) return false;
+
+        n.x += t.x;
+        n.y += t.y;
+        // n.z += t.z; - VLFeat is not changing levels !!!
+
+        return true;
+    }
+};
+
+template<>
+class ModeFunctions<Config::OpenCV>
+{
+public:
+    inline __device__
+    bool first_contrast_ok( const float val ) const
+    {
+        return ( fabs( val ) >= d_threshold );
+    }
+
+    inline __device__
+    bool refine( const float3& d, int3& n, const int width, const int height, const int maxlevel )
+    {
+        float3 t;
+
+        t.x = ((d.x >= 0.5f && n.x < width-2) ?  1 : 0 )
+            + ((d.x <= -0.5f && n.x > 1)? -1 : 0 );
+
+        t.y = ((d.y >= 0.5f && n.y < height-2)  ?  1 : 0 )
+            + ((d.y <= -0.5f && n.y > 1) ? -1 : 0 );
+
+        t.z = ((d.z >= 0.5f && n.z < maxlevel-1)  ?  1 : 0 )
+            + ((d.z <= -0.5f && n.z > 1) ? -1 : 0 );
+
+        if( t.x == 0 && t.y == 0 && t.z == 0 ) return false;
+
+        n.x += t.x;
+        n.y += t.y;
+        n.z += t.z;
+
+        return true;
+    }
+};
+
+template<>
+class ModeFunctions<Config::PopSift>
+{
+public:
+    inline __device__
+    bool first_contrast_ok( const float val ) const
+    {
+        return ( fabs( val ) >= 1.6 * d_threshold );
+    }
+
+    inline __device__
+    bool refine( float3& d, int3& n, int width, int height, int maxlevel )
+    {
+        float3 t;
+
+        t.x = ((d.x >= 0.6f && n.x < width-2) ?  1 : 0 )
+            + ((d.x <= -0.6f && n.x > 1)? -1 : 0 );
+
+        t.y = ((d.y >= 0.6f && n.y < height-2)  ?  1 : 0 )
+            + ((d.y <= -0.6f && n.y > 1) ? -1 : 0 );
+
+        t.z = ((d.z >= 0.6f && n.z < maxlevel-1)  ?  1 : 0 )
+            + ((d.z <= -0.6f && n.z > 1) ? -1 : 0 );
+
+        if( t.x == 0 && t.y == 0 && t.z == 0 ) return false;
+
+        n.x += t.x;
+        n.y += t.y;
+        n.z += t.z;
+
+        return true;
+    }
+};
+
+template<int sift_mode>
 __device__
 bool find_extrema_in_dog_v6_sub( cudaTextureObject_t dog,
                                  int                 debug_octave,

    /* reject condition: tr(H)^2/det(H) < (r+1)^2/r */
    if( edgeval >= (d_edge_limit+1.0f)*(d_edge_limit+1.0f)/d_edge_limit ) {
#ifdef PRINT_EXTREMA_DEBUG_INFO
        printf("Found an extremum at %d %d (o=%d) - edge tresh failed\n", x+1, y+1, debug_octave );
#endif // PRINT_EXTREMA_DEBUG_INFO
        return false;
    }

    ec.sigma   = d_sigma0 * pow(d_sigma_k, sn); // * 2;
        // const float sigma_k = powf(2.0f, 1.0f / levels );
#ifdef PRINT_EXTREMA_DEBUG_INFO
    printf("Found an extremum at %d %d (o=%d,l=%d)     -> x:%.1f y:%.1f z:%.1f\n", x+1, y+1, debug_octave, level, xn, yn, sn );
#endif // PRINT_EXTREMA_DEBUG_INFO

    ec.orientation = 0;
