 __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 );
-    assert( getWidth()  == _d_map.cols );
-    assert( getHeight() == _d_map.rows );
-    assert( getWidth()  == _d_hyst_edges.cols );
-    assert( getHeight() == _d_hyst_edges.rows );
-
-#ifndef NDEBUG
+
     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>>>
-        ( block, grid, _d_hysteresis_block_counter, _d_hyst_edges, _d_map );
+        ( 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;

}

__device__
void load( cv::cuda::PtrStepSzb img )
{
#if 0
    //
    // After several dead ends, this is the most trivial, most
    // brute-force approach for loading that I could find.
    //
    const int srcidx = blockIdx.x * HYST_W + threadIdx.x;
    const int srcidy = blockIdx.y * HYST_H + threadIdx.y;
    // const int dstidx = threadIdx.x;
    // const int dstidy = threadIdx.y;

    volatile uint8_t val[3][3];
    val[0][0] = get( img, srcidx-1, srcidy-1 );
    val[0][1] = get( img, srcidx  , srcidy-1 );
    val[0][2] = get( img, srcidx+1, srcidy-1 );
    val[1][0] = get( img, srcidx-1, srcidy   );
    val[1][1] = get( img, srcidx  , srcidy   );
    val[1][2] = get( img, srcidx+1, srcidy   );
    val[2][0] = get( img, srcidx-1, srcidy+1 );
    val[2][1] = get( img, srcidx  , srcidy+1 );
    val[2][2] = get( img, srcidx+1, srcidy+1 );

    assert( val[0][0] <= 2 ); assert( val[0][1] <= 2 ); assert( val[0][2] <= 2 ); assert( val[1][0] <= 2 ); assert( val[1][1] <= 2 ); assert( val[1][2] <= 2 ); assert( val[2][0] <= 2 ); assert( val[2][1] <= 2 ); assert( val[2][2] <= 2 );

    array[threadIdx.y  ][threadIdx.x  ] = val[0][0];
    array[threadIdx.y  ][threadIdx.x+2] = val[0][2];
    if( threadIdx.y >= HYST_H - 2 ) {
        array[threadIdx.y+2][threadIdx.x  ] = val[2][0];
        array[threadIdx.y+2][threadIdx.x+2] = val[2][2];
    }
    __syncthreads();
#else
    const int srcidx = blockIdx.x * HYST_W + threadIdx.x;
    const int srcidy = blockIdx.y * HYST_H + threadIdx.y;

    uint8_t val_0_0;
    val_0_0 = get( img, srcidx-1, srcidy-1 );
    array[threadIdx.y  ][threadIdx.x  ] = val_0_0;

    if( threadIdx.x >= HYST_W - 2 ) {
        uint8_t val_0_2;
        val_0_2 = get( img, srcidx+1, srcidy-1 );
        array[threadIdx.y  ][threadIdx.x+2] = val_0_2;
    }
    if( threadIdx.y >= HYST_H - 2 ) {
        uint8_t val_2_0;
        val_2_0 = get( img, srcidx-1, srcidy+1 );
        array[threadIdx.y+2][threadIdx.x  ] = val_2_0;
        if( threadIdx.x >= HYST_W - 2 ) {
            uint8_t val_2_2;
            val_2_2 = get( img, srcidx+1, srcidy+1 );
            array[threadIdx.y+2][threadIdx.x+2] = val_2_2;
        }
    }
    __syncthreads();
#endif
}

__device__
void store( cv::cuda::PtrStepSzb img )
{
    const int dstidx  = blockIdx.x * HYST_W + threadIdx.x;
    const int dstidy  = blockIdx.y * HYST_H + threadIdx.y;

    volatile uint8_t val;
    val = array[threadIdx.y+1][threadIdx.x+1];
    assert( val <= 2 );
    // if( outOfBounds( dstidx, dstidy, img ) ) return;
    if( dstidx < img.cols && dstidy < img.rows ) {
        img.ptr(dstidy)[dstidx] =  val;
    }
    // __syncthreads();
}

__device__
inline
bool update_edge_pixel( )
{
    uint8_t val[3][3];
    val[0][0] = array[threadIdx.y  ][threadIdx.x  ];
    val[0][1] = array[threadIdx.y  ][threadIdx.x+1];
    val[0][2] = array[threadIdx.y  ][threadIdx.x+2];
    val[1][0] = array[threadIdx.y+1][threadIdx.x  ];
    val[1][1] = array[threadIdx.y+1][threadIdx.x+1];
    val[1][2] = array[threadIdx.y+1][threadIdx.x+2];
    val[2][0] = array[threadIdx.y+2][threadIdx.x  ];
    val[2][1] = array[threadIdx.y+2][threadIdx.x+1];
    val[2][2] = array[threadIdx.y+2][threadIdx.x+2];

    assert( val[0][0] <= 2 );
    assert( val[0][1] <= 2 );

        if( nothing_changed ) {
            atomicSub( block_counter, 1 );
        }
    }
    __syncthreads();
    return nothing_changed;
