 inline
 uint32_t get( cv::cuda::PtrStepSz32u img, const int idx, const int idy )
 {
-#if 1
-    return img.ptr( clamp( idy, img.rows ) )[ clamp( idx, img.cols ) ];
-#else
-    int x = clamp( idx, img.cols );
-    int y = clamp( idy, img.rows );
-    assert( x >= 0 );
-    assert( y >= 0 );
-    assert( x < img.cols );
-    assert( y < img.rows );
-    uint8_t  val = img.ptr(y)[x];
-    if( val > 2 ) {
-        printf("idx=%d -> x=%d, idy=%d -> y=%d, img.cols=%d img.rows=%d val=%d\n",
-            idx, x, idy, y, img.cols, img.rows, val );
-        assert( val <= 2 );
-    }
-    return val;
-#endif
+    if( idx < 0 || idy < 0 || idx >= img.cols || idy >= img.rows ) {
+        return 0;
+    }
+    return img.ptr(idy)[idx];
 }
 
 __device__

            grid.x, grid.y, grid.z, block.x, block.y, block.z );

    bool first_time = true;
    int loopcount = 0;
    do
    {
        *block_counter = 0;

}

__global__
void edge_first( cv::cuda::PtrStepSzb img, int* block_counter, cv::cuda::PtrStepSzb src, int debug_roundcount )
{
    // const int idx  = blockIdx.x * HYST_W + threadIdx.x;
    // const int idy  = blockIdx.y * HYST_H + threadIdx.y;
