     // send data to device
     gpuErrchk(cudaMemcpy(g_src_data, (*src_bitmap)->data, g_data_size, cudaMemcpyHostToDevice));
 
-    // uint8_t are_identical_bitmaps = 0;
+    uint8_t are_identical_bitmaps = 0;
     int iterations = 0;
     do {
         int skeletonize_pass_shared_mem_size = (block_dim.x + PAD_LEFT + PAD_RIGHT) * (1 + PAD_TOP + PAD_BOTTOM) * sizeof(uint8_t);

#define PAD_BOTTOM 1
#define PAD_RIGHT 1

void and_reduction(uint8_t* g_data, int g_size, dim3 grid_dim, dim3 block_dim) {
    int shared_mem_size = block_dim.x * sizeof(uint8_t);

    // iterative reductions of g_data
    do {
        and_reduction<<<grid_dim, block_dim, shared_mem_size>>>(g_data, g_size);
        gpuErrchk(cudaPeekAtLastError());
        gpuErrchk(cudaDeviceSynchronize());


        gpuErrchk(cudaPeekAtLastError());
        gpuErrchk(cudaDeviceSynchronize());

        and_reduction(g_src_data, g_dst_data, g_equ_data, (*src_bitmap)->width * (*src_bitmap)->height, grid_dim, block_dim);

        // bring reduced bitmap equality information back from device
        gpuErrchk(cudaMemcpy(&are_identical_bitmaps, g_equ_data, 1 * sizeof(uint8_t), cudaMemcpyDeviceToHost));

        swap_bitmaps((void**) &g_src_data, (void**) &g_dst_data);

        iterations++;
        printf(".");
        fflush(stdout);
    } while (!are_identical_bitmaps);

    // bring dst_bitmap back from device
    gpuErrchk(cudaMemcpy((*dst_bitmap)->data, g_dst_data, g_data_size, cudaMemcpyDeviceToHost));

    // free memory on device
    gpuErrchk(cudaFree(g_src_data));
