     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;

}

__device__
void load( cv::cuda::PtrStepSzb img )
{
#if 0
    //
    // After several dead ends, this is the most trivial, most
    // brute-force approach for loading that I could find.
    //
    const int srcidx = blockIdx.x * HYST_W + threadIdx.x;
    const int srcidy = blockIdx.y * HYST_H + threadIdx.y;
    // const int dstidx = threadIdx.x;
    // const int dstidy = threadIdx.y;

    volatile uint8_t val[3][3];
    val[0][0] = get( img, srcidx-1, srcidy-1 );
    val[0][1] = get( img, srcidx  , srcidy-1 );
    val[0][2] = get( img, srcidx+1, srcidy-1 );
    val[1][0] = get( img, srcidx-1, srcidy   );
    val[1][1] = get( img, srcidx  , srcidy   );
    val[1][2] = get( img, srcidx+1, srcidy   );
    val[2][0] = get( img, srcidx-1, srcidy+1 );
    val[2][1] = get( img, srcidx  , srcidy+1 );
    val[2][2] = get( img, srcidx+1, srcidy+1 );

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

    array[threadIdx.y  ][threadIdx.x  ] = val[0][0];
    array[threadIdx.y  ][threadIdx.x+2] = val[0][2];
    if( threadIdx.y >= HYST_H - 2 ) {
        array[threadIdx.y+2][threadIdx.x  ] = val[2][0];
        array[threadIdx.y+2][threadIdx.x+2] = val[2][2];
    }
    __syncthreads();
#else
    const int srcidx = blockIdx.x * HYST_W + threadIdx.x;
    const int srcidy = blockIdx.y * HYST_H + threadIdx.y;

    uint8_t val_0_0;
    val_0_0 = get( img, srcidx-1, srcidy-1 );
    array[threadIdx.y  ][threadIdx.x  ] = val_0_0;

    if( threadIdx.x >= HYST_W - 2 ) {
        uint8_t val_0_2;
        val_0_2 = get( img, srcidx+1, srcidy-1 );
        array[threadIdx.y  ][threadIdx.x+2] = val_0_2;
    }
    if( threadIdx.y >= HYST_H - 2 ) {
        uint8_t val_2_0;
        val_2_0 = get( img, srcidx-1, srcidy+1 );
        array[threadIdx.y+2][threadIdx.x  ] = val_2_0;
        if( threadIdx.x >= HYST_W - 2 ) {
            uint8_t val_2_2;
            val_2_2 = get( img, srcidx+1, srcidy+1 );
            array[threadIdx.y+2][threadIdx.x+2] = val_2_2;
        }
    }
    __syncthreads();
#endif
}

__device__
void store( cv::cuda::PtrStepSzb img )
{
    const int dstidx  = blockIdx.x * HYST_W + threadIdx.x;
    const int dstidy  = blockIdx.y * HYST_H + threadIdx.y;

    volatile uint8_t val;
    val = array[threadIdx.y+1][threadIdx.x+1];
    assert( val <= 2 );
    // if( outOfBounds( dstidx, dstidy, img ) ) return;
    if( dstidx < img.cols && dstidy < img.rows ) {
        img.ptr(dstidy)[dstidx] =  val;
    }
    // __syncthreads();
}

__device__
inline
bool update_edge_pixel( )
{
    uint8_t val[3][3];
    val[0][0] = array[threadIdx.y  ][threadIdx.x  ];
    val[0][1] = array[threadIdx.y  ][threadIdx.x+1];
    val[0][2] = array[threadIdx.y  ][threadIdx.x+2];
    val[1][0] = array[threadIdx.y+1][threadIdx.x  ];
    val[1][1] = array[threadIdx.y+1][threadIdx.x+1];
    val[1][2] = array[threadIdx.y+1][threadIdx.x+2];
    val[2][0] = array[threadIdx.y+2][threadIdx.x  ];
    val[2][1] = array[threadIdx.y+2][threadIdx.x+1];
    val[2][2] = array[threadIdx.y+2][threadIdx.x+2];

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


namespace hysteresis
{
#define HYST_H   32
#define HYST_W   32

