                                          &block_counter,
                                          sizeof(int), _stream );
         if( first_time ) {
+            cerr << "1" << endl;
             hysteresis::edge_first
                 <<<grid,block,0,_stream>>>
                 ( _d_hyst_edges,

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

    uint32_t* reinterpret_array = reinterpret_cast<uint32_t*>(array);

    volatile uint32_t val;
    val = reinterpret_array[threadIdx.y+1][threadIdx.x+1];

    if( dstidx < img.cols && dstidy < img.rows ) {
        img.ptr(dstidy)[dstidx] =  val;

                  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

__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 );
