     input.cols = img.cols / 4;
     load( input );
 
-    bool something_changed = edge( block_counter );
+    bool something_changed = edge( block_counter, debug_roundcount );
 
     if( __any( something_changed ) ) {
-        store( input );
+        store( input, false );
     }
 }
 

                <<<grid,block>>>
                ( img,
                  block_counter,
                  src );
            first_time = false;
        } else {
            hysteresis::edge_second
                <<<grid,block>>>
                ( img,
                  block_counter );
        }
        cudaDeviceSynchronize( );
        printf( "block_counter=%d\n", *block_counter);
        assert( *block_counter <= grid.x * grid.y );
    }
    while( loopcount++ < 30 ); // *block_counter > 0 );
    printf( "Leave %s\n", __FUNCTION__ );
}
#endif // USE_SEPARABLE_COMPILATION

        if( again ) something_changed = true;

        /* this should not be necessary ... */
        debug_inner_loop_count++;
    }

    return something_changed;
}

__device__
bool edge( int* block_counter, int debug_roundcount )
{
    bool something_changed = edge_block_loop( debug_roundcount );
    if( threadIdx.x == 0 && threadIdx.y == 0 ) {
        if( something_changed ) {
            if( debug_roundcount > 25 ) {
                printf("Something changed in block (%d,%d)\n", blockIdx.x, blockIdx.y );
            }
            atomicAdd( block_counter, 1 );
        }
    }
