     assert( val[2][1] <= 2 );
     assert( val[2][2] <= 2 );
 
-    bool    change = false;
+    bool inc = false;
+    bool dec = false;
 
     if( val[1][1] == 1 ) {
-        change = ( 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 );
-        val[1][1] = change ? 2 : 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[threadIdx.y+1][threadIdx.x+1] = val[1][1];
-
-#if 0
-    uint8_t test = array[threadIdx.y+1][threadIdx.x+1];
-    assert( test <= 2 );
-    assert( not change || test == 2 );
-#endif
-
-    return change;
+    array[y+1][x+1] = val[1][1];
+
+    return ( inc || dec );
 }
 
 __device__
 bool edge_block_loop( )
 {
-    __shared__ bool continuation[HYST_H];
+    __shared__ volatile bool continuation[HYST_H];
     bool            again = true;
     bool            nothing_changed = true;
     bool            line_changed = false;


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


    // if( outOfBounds( idx, idy, img ) ) return;
    // uint8_t val = src.ptr(idy)[idx];
    // img.ptr(idy)[idx] = val;
    PtrStepSz32u input( src );
    input.cols = src.cols / 4;
    load( src );

    edge( block_counter );

    PtrStepSz32u output( img );
    output.cols = img.cols / 4;
    store( img );
}

__global__
void edge_second( cv::cuda::PtrStepSzb img, int* block_counter )
{
    PtrStepSz32u input( src );
    input.cols = src.cols / 4;
    load( img );

    bool nothing_changed = edge( block_counter );
    if( not nothing_changed ) {
        store( img );
