         dec = ( val[0][0] == 0 && val[0][1] == 0 && val[0][2] == 0 &&
                 val[1][0] == 0 &&                   val[1][2] == 0 &&
                 val[2][0] == 0 && val[2][1] == 0 && val[2][2] == 0 );
-        val[1][1] = inc ? 2 : dec ? 0 : 1 ;
+        debug_v = val[1][1] = inc ? 2 : dec ? 0 : 1 ;
     }
     __syncthreads();
+    assert( debug_v == val[1][1] );
     array[y+1][x+1] = val[1][1];
 
     return ( inc || dec );

        /* 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__

    cudaEvent_t before_hyst, after_hyst;
    float ms;

    cudaEventCreate( &before_hyst );
    cudaEventCreate( &after_hyst );
    cudaEventRecord( before_hyst, _stream );
    cerr << "0" << endl;
    hyst_outer_loop
        <<<1,1,0,_stream>>>
        ( getWidth(), getHeight(), _d_hysteresis_block_counter, _d_hyst_edges, _d_map );
    cudaEventRecord( after_hyst, _stream );
    cerr << "0.1" << endl;
    cudaEventSynchronize( after_hyst );
    cerr << "0.2" << endl;
    cudaEventElapsedTime( &ms, before_hyst, after_hyst );
    cudaEventDestroy( before_hyst );
    cudaEventDestroy( after_hyst );
    std::cerr << "Hyst took " << ms << " ms" << std::endl;
#else // USE_SEPARABLE_COMPILATION
    bool first_time = true;
