     cudaEvent_t before_hyst, after_hyst;
     float ms;
 
-    cudaCreateEvent( &before_hyst );
-    cudaCreateEvent( &after_hyst );
+    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 );
-    cudaEventDestory( before_hyst );
-    cudaEventDestory( after_hyst );
+    cudaEventDestroy( before_hyst );
+    cudaEventDestroy( after_hyst );
     std::cerr << "Hyst took " << ms << " ms" << std::endl;
 #else // USE_SEPARABLE_COMPILATION
     bool first_time = true;

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