 }
 
 __device__ uint8_t block_and_reduce(uint8_t* s_data) {
-    int tid = threadIdx.y * blockDim.x + threadIdx.x;
-
-    for (int s = ((blockDim.x * blockDim.y) / 2); s > 0; s >>= 1) {
-        if (tid < s) {
-            s_data[tid] &= s_data[tid + s];
+    for (int s = (blockDim.x / 2); s > 0; s >>= 1) {
+        if (threadIdx.x < s) {
+            s_data[threadIdx.x] &= s_data[threadIdx.x + s];
         }
         __syncthreads();
     }

#define PAD_BOTTOM 1
#define PAD_RIGHT 1

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

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

        g_width = grid_dim.x;
        g_height = grid_dim.y;
        grid_dim.x = ceil(grid_dim.x / ((double) block_dim.x));
        grid_dim.y = ceil(grid_dim.y / ((double) block_dim.y));
    } while ((g_width * g_height) != 1);
}

// Adapted for 2D arrays from Nvidia cuda SDK samples
__global__ void and_reduction(uint8_t* g_data, int g_width, int g_height) {
    // shared memory for tile
    extern __shared__ uint8_t s_data[];

    int g_row = blockIdx.y * blockDim.y + threadIdx.y;
    int g_col = blockIdx.x * blockDim.x + threadIdx.x;

    int tid = threadIdx.y * blockDim.x + threadIdx.x;

    // Load equality values into shared memory tile. We use 1 as the default
    // value, as it is an AND reduction
    s_data[tid] = is_outside_image(g_row, g_col, g_width, g_height) ? 1 : global_mem_read(g_data, g_row, g_col, g_width, g_height);
    __syncthreads();

    // do reduction in shared memory
    uint8_t write_data = block_and_reduce(s_data);

    // write result for this block to global memory
    if (tid == 0) {
        global_mem_write(g_data, blockIdx.y, blockIdx.x, gridDim.x, gridDim.y, write_data);
    }
}


    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);
        skeletonize_pass<<<grid_dim, block_dim, skeletonize_pass_shared_mem_size>>>(g_src_data, g_dst_data, (*src_bitmap)->width, (*src_bitmap)->height);
        gpuErrchk(cudaPeekAtLastError());
        gpuErrchk(cudaDeviceSynchronize());

        pixel_equality<<<grid_dim, block_dim>>>(g_src_data, g_dst_data, g_equ_data, (*src_bitmap)->width * (*src_bitmap)->height);
        gpuErrchk(cudaPeekAtLastError());
        gpuErrchk(cudaDeviceSynchronize());

        and_reduction(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));
