     int iterations = 0;
     do {
         gpuErrchk(cudaMemcpy(g_dst_data, g_src_data, g_data_size, cudaMemcpyDeviceToDevice));
+        gpuErrchk(cudaMemset(g_equ_data, 1, g_data_size));
 
         int skeletonize_pass_s_src_size = (block_dim.x + PAD_LEFT + PAD_RIGHT) * (block_dim.y + PAD_TOP + PAD_BOTTOM) * sizeof(uint8_t);
         int skeletonize_pass_s_equ_size = block_dim.x * block_dim.y * sizeof(uint8_t);

        uint8_t thinning_cond_4 = (((P2 & P4 & P6) == 0) | (TR_P4 != 1));
        uint8_t thinning_cond_ok = thinning_cond_1 & thinning_cond_2 & thinning_cond_3 & thinning_cond_4;

        g_dst_next = BINARY_WHITE + ((1 - thinning_cond_ok) * s_src[s_src_row * s_src_width + s_src_col]);

        // write g_dst_next to g_dst
        global_mem_write(g_dst, g_row, g_col, g_width, g_height, g_dst_next);
    }

    // write pixel equality information to g_equ
    uint8_t write_data = (s_src[(s_src_row) * s_src_width + (s_src_col)] == g_dst_next);
    global_mem_write(g_equ, g_row, g_col, g_width, g_height, write_data);
}

// Computes the number of white to black transitions around a pixel.

    int s_equ_row = threadIdx.y;
    int s_equ_col = threadIdx.x;
    int s_equ_width = blockDim.x;

    // load g_src into shared memory
    load_s_src(g_src, g_row, g_col, g_width, g_height, s_src, s_src_row, s_src_col, s_src_width);
