     // if( outOfBounds( idx, idy, img ) ) return;
     // uint8_t val = src.ptr(idy)[idx];
     // img.ptr(idy)[idx] = val;
+    PtrStepSz32u input( src );
+    input.cols = src.cols / 4;
     load( src );
-    bool nothing_changed = edge( block_counter );
+
+    edge( block_counter );
+
+    PtrStepSz32u output( img );
+    output.cols = img.cols / 4;
     store( img );
 }
 
 __global__
 void edge_second( cv::cuda::PtrStepSzb img, int* block_counter )
 {
+    PtrStepSz32u input( src );
+    input.cols = src.cols / 4;
     load( img );
+
     bool nothing_changed = edge( block_counter );
     if( not nothing_changed ) {
         store( img );

#error The code requires W<=32 and H<=W
#endif

__shared__ volatile uint8_t array[HYST_H+2][HYST_W+2];

__device__
inline
uint8_t get( cv::cuda::PtrStepSzb img, const int idx, const int idy )
{
#if 1
    return img.ptr( clamp( idy, img.rows ) )[ clamp( idx, img.cols ) ];

__host__
void Frame::applyHyst( const cctag::Parameters & params )
{
    assert( getWidth()  == _d_map.cols );
    assert( getHeight() == _d_map.rows );
    assert( getWidth()  == _d_hyst_edges.cols );
    assert( getHeight() == _d_hyst_edges.rows );

#ifndef NDEBUG
    dim3 block;
    dim3 grid;
    block.x = HYST_W;
    block.y = HYST_H;
    grid.x  = grid_divide( getWidth(),   HYST_W );
    grid.y  = grid_divide( getHeight(),  HYST_H );

    verify_map_valid
        <<<grid,block,0,_stream>>>
        ( _d_map, _d_hyst_edges, getWidth(), getHeight() );
#endif

#if defined(USE_SEPARABLE_COMPILATION)
    cudaEvent_t before_hyst, after_hyst;
    float ms;

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