 
 namespace hysteresis
 {
-// HYST    4 x 32 -> >10
-// HYST    2 x 32 -> >11.5
-// HYST   16 x 16 -> 9.380 - 10.426
-// HYST    8 x 32 -> 9.6 - 9.9
-// HYST   32 x 32 -> 9.460 - 9.793
-// HYST    8 x  8 -> 9.7 - 10.16
 #define HYST_H   32
 #define HYST_W   32
 


#if defined(USE_SEPARABLE_COMPILATION)
__global__
void hyst_outer_loop( dim3 block, dim3 grid, int* block_counter, cv::cuda::PtrStepSzb img, cv::cuda::PtrStepSzb src )
{
    bool first_time = true;
    int gridsize = grid.x * grid.y;
    do

}

__device__
void load( cv::cuda::PtrStepSz32u img )
{
    const int srcidx = blockIdx.x * HYST_W + threadIdx.x;
    const int srcidy = blockIdx.y * HYST_H + threadIdx.y;

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

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

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

__device__
void store( cv::cuda::PtrStepSz32u img )
{
    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;
    }
}

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

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