     output.step = img.step;
     output.rows = img.rows;
     output.cols = img.cols / 4;
-    store( output );
+    store( output, false );
 }
 
 __global__
-void edge_second( cv::cuda::PtrStepSzb img, int* block_counter )
+void edge_second( cv::cuda::PtrStepSzb img, int* block_counter, int debug_roundcount )
 {
     cv::cuda::PtrStepSz32u input;
     input.data = reinterpret_cast<uint32_t*>(img.data);

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
}

__device__

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

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