 
     while( again ) {
         assert( ct <= HYST_W*HYST_H );
-        bool mark    = update_edge_pixel( );
-        __threadfence();
+        bool mark = false;
+        mark = mark || update_edge_pixel( threadIdx.y, x    );
+        mark = mark || update_edge_pixel( threadIdx.y, x+32 );
+        mark = mark || update_edge_pixel( threadIdx.y, x+64 );
+        mark = mark || update_edge_pixel( threadIdx.y, x+96 );
+
+        /* make sure all updated pixel are written back to
+         * shared memory before continuation[] is modified */
+        __threadfence_block();
+
+        /* every row checks whether any pixel has been changed */
         line_changed = __any( mark );
-        if( threadIdx.x == 0 ) continuation[threadIdx.y] = line_changed;
+
+        /* the first thread of each row write the result to continuation[] */
+        if( threadIdx.x == 0 ) continuation[y] = line_changed;
+
+        /* wait for all rows to fulfill the operation (and to assure that
+         * results in continuation[] are visible to all threads, because
+         * threadfence() is implied by syncthreads() */
         __syncthreads();
+
+        /* Each thread in a warp reads continuation for one row.
+         * Redundant, but I have no better idea for spreading the result
+         * to all warps. */
         mark = threadIdx.x < HYST_H ? continuation[threadIdx.x] : false;
+
+        /* Finally, all 32x32 threads know whether at least one of them
+         * has changed a pixel.
+         * If there has been any change in this round, try to spread
+         * the change further.
+         */
         again = __any( mark );
+
+        /* 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++;
     }
 


#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


namespace hysteresis
{
#define HYST_H   32
#define HYST_W   32

