                 <<<grid,block>>>
                 ( img,
                   block_counter,
-                  src,
-                  debug_roundcount );
+                  src );
             first_time = false;
         } else {
             hysteresis::edge_second
                 <<<grid,block>>>
                 ( img,
-                  block_counter,
-                  debug_roundcount );
+                  block_counter );
         }
         cudaDeviceSynchronize( );
-        printf( "width=%d height=%d block_counter=%d\n", width, height, *block_counter);
         assert( *block_counter <= grid.x * grid.y );
     }
-    while( *block_counter > 0 && debug_roundcount++ < 30 ); // *block_counter > 0 );
-    printf( "Leave %s\n", __FUNCTION__ );
+    while( *block_counter > 0 );
 }
 #endif // USE_SEPARABLE_COMPILATION
 
 __host__
 void Frame::applyHyst( const cctag::Parameters & params )
 {
-    cerr << "Enter " << __FUNCTION__ << endl;
     assert( getWidth()  == _d_map.cols );
     assert( getHeight() == _d_map.rows );
     assert( getWidth()  == _d_hyst_edges.cols );

    const int dstidx  = blockIdx.x * HYST_W + threadIdx.x;
    const int dstidy  = blockIdx.y * HYST_H + threadIdx.y;

#if 1
    if( dstidx*sizeof(uint32_t) < img.step && dstidy < img.rows ) {
        // volatile uint32_t* shared_line = reinterpret_cast<volatile uint32_t*>(array[threadIdx.y+1]);
        volatile uint32_t* shared_line = reinterpret_cast<volatile uint32_t*>(&array[threadIdx.y+1][0]);

}

__device__
bool edge( int* block_counter )
{
    bool something_changed = edge_block_loop( );
    if( threadIdx.x == 0 && threadIdx.y == 0 ) {
        if( something_changed ) {
            atomicAdd( block_counter, 1 );
        }
    }
