 
 namespace hysteresis
 {
-// HYST    4 x 32 -> >10
-// HYST    2 x 32 -> >11.5
-// HYST   16 x 16 -> 9.380 - 10.426
-// HYST    8 x 32 -> 9.6 - 9.9
-// HYST   32 x 32 -> 9.460 - 9.793
-// HYST    8 x  8 -> 9.7 - 10.16
 #define HYST_H   32
 #define HYST_W   32
 

#error The code requires W<=32 and H<=W
#endif

__shared__ volatile uint8_t array[HYST_H+2][HYST_W+2];

__device__
inline
uint8_t get( cv::cuda::PtrStepSzb img, const int idx, const int idy )
{
#if 1
    return img.ptr( clamp( idy, img.rows ) )[ clamp( idx, img.cols ) ];

    assert( val[2][1] <= 2 );
    assert( val[2][2] <= 2 );

    bool inc = false;
    bool dec = false;

    if( val[1][1] == 1 ) {
        inc = ( val[0][0] == 2 || val[0][1] == 2 || val[0][2] == 2 ||
                val[1][0] == 2 ||                   val[1][2] == 2 ||
                val[2][0] == 2 || val[2][1] == 2 || val[2][2] == 2 );
        dec = ( val[0][0] == 0 && val[0][1] == 0 && val[0][2] == 0 &&
                val[1][0] == 0 &&                   val[1][2] == 0 &&
                val[2][0] == 0 && val[2][1] == 0 && val[2][2] == 0 );
        val[1][1] = inc ? 2 : dec ? 0 : 1 ;
    }
    __syncthreads();
    array[y+1][x+1] = val[1][1];

    return ( inc || dec );
}

__device__
bool edge_block_loop( )
{
    __shared__ volatile bool continuation[HYST_H];
    bool            again = true;
    bool            nothing_changed = true;
    bool            line_changed = false;
