                   block_counter );
         }
         cudaDeviceSynchronize( );
-    }
-    while( *block_counter > 0 );
+        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
 
 __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 );

        /* Every threads needs to know whether any pixel was changed in
         * any round of the loop because egde_second() uses this return
         * value to write back to global memory using a different alignment. */
        if( again ) nothing_changed = false;

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

    return nothing_changed;
}

__device__
bool edge( int* block_counter )
{
    bool nothing_changed = edge_block_loop( );
    if( threadIdx.x == 0 && threadIdx.y == 0 ) {
        if( nothing_changed ) {
            atomicSub( block_counter, 1 );
        }
    }
    __syncthreads();
    return nothing_changed;
}

__global__

                  _d_map );
            first_time = false;
        } else {
            cerr << "2" << endl;
            hysteresis::edge_second
                <<<grid,block,0,_stream>>>
                ( _d_hyst_edges,
