     uint8_t are_identical_bitmaps = 0;
     int iterations = 0;
     do {
-        int skeletonize_pass_shared_mem_size = (block_dim.x + PAD_LEFT + PAD_RIGHT) * (block_dim.y + PAD_TOP + PAD_BOTTOM) * sizeof(uint8_t);
+        int skeletonize_pass_shared_mem_size = (((block_dim.x + PAD_LEFT + PAD_RIGHT) * (block_dim.y + PAD_TOP + PAD_BOTTOM)) + (block_dim.x * block_dim.y)) * 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());

    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;

    uint8_t write_data = BINARY_WHITE + ((1 - thinning_cond_ok) * s_src[s_row * s_width + s_col]);
    global_mem_write(g_dst, g_row, g_col, g_width, g_height, write_data);
}


// Performs 1 iteration of the thinning algorithm.
__global__ void skeletonize_pass(uint8_t* g_src, uint8_t* g_dst, int g_width, int g_height) {
    // shared memory for tile
    extern __shared__ uint8_t s_data[];

    uint8_t* s_src = &s_data[0];
    uint8_t* s_dst = &s_data[(blockDim.x + PAD_LEFT + PAD_RIGHT) * (blockDim.y + PAD_TOP + PAD_BOTTOM)];

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

    int s_src_row = threadIdx.y + PAD_TOP;
    int s_src_col = threadIdx.x + PAD_LEFT;
    int s_src_width = blockDim.x + PAD_LEFT + PAD_RIGHT;

    int s_dst_row = threadIdx.y;
    int s_dst_col = threadIdx.x;
    int s_dst_width = blockDim.x;

    // load g_src into shared memory
    if ((threadIdx.y == 0) & (threadIdx.x == 0)) {
        // top-left corner
        s_src[(s_src_row - 2) * s_src_width + (s_src_col - 2)] = global_mem_read(g_src, g_row - 2, g_col - 2, g_width, g_height);
        s_src[(s_src_row - 2) * s_src_width + (s_src_col - 1)] = global_mem_read(g_src, g_row - 2, g_col - 1, g_width, g_height);
        s_src[(s_src_row - 2) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row - 2, g_col, g_width, g_height);

        s_src[(s_src_row - 1) * s_src_width + (s_src_col - 2)] = global_mem_read(g_src, g_row - 1, g_col - 2, g_width, g_height);
        s_src[(s_src_row - 1) * s_src_width + (s_src_col - 1)] = global_mem_read(g_src, g_row - 1, g_col - 1, g_width, g_height);
        s_src[(s_src_row - 1) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row - 1, g_col, g_width, g_height);

        s_src[(s_src_row) * s_src_width + (s_src_col - 2)] = global_mem_read(g_src, g_row, g_col - 2, g_width, g_height);
        s_src[(s_src_row) * s_src_width + (s_src_col - 1)] = global_mem_read(g_src, g_row, g_col - 1, g_width, g_height);
        s_src[(s_src_row) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row, g_col, g_width, g_height);
    } else if ((threadIdx.y == (blockDim.y - 1)) & (threadIdx.x == 0)) {
        // bottom-left corner
        s_src[(s_src_row) * s_src_width + (s_src_col - 2)] = global_mem_read(g_src, g_row, g_col - 2, g_width, g_height);
        s_src[(s_src_row) * s_src_width + (s_src_col - 1)] = global_mem_read(g_src, g_row, g_col - 1, g_width, g_height);
        s_src[(s_src_row) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row, g_col, g_width, g_height);

        s_src[(s_src_row + 1) * s_src_width + (s_src_col - 2)] = global_mem_read(g_src, g_row + 1, g_col - 2, g_width, g_height);
        s_src[(s_src_row + 1) * s_src_width + (s_src_col - 1)] = global_mem_read(g_src, g_row + 1, g_col - 1, g_width, g_height);
        s_src[(s_src_row + 1) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row + 1, g_col, g_width, g_height);
    } else if ((threadIdx.y == (blockDim.y - 1)) & (threadIdx.x == (blockDim.x - 1))) {
        // bottom-right corner
        s_src[(s_src_row) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row, g_col, g_width, g_height);
        s_src[(s_src_row) * s_src_width + (s_src_col + 1)] = global_mem_read(g_src, g_row, g_col + 1, g_width, g_height);

        s_src[(s_src_row + 1) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row + 1, g_col, g_width, g_height);
        s_src[(s_src_row + 1) * s_src_width + (s_src_col + 1)] = global_mem_read(g_src, g_row + 1, g_col + 1, g_width, g_height);
    } else if ((threadIdx.y == 0) & (threadIdx.x == (blockDim.x - 1))) {
        // top-right corner
        s_src[(s_src_row - 2) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row - 2, g_col, g_width, g_height);
        s_src[(s_src_row - 2) * s_src_width + (s_src_col + 1)] = global_mem_read(g_src, g_row - 2, g_col + 1, g_width, g_height);

        s_src[(s_src_row - 1) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row - 1, g_col, g_width, g_height);
        s_src[(s_src_row - 1) * s_src_width + (s_src_col + 1)] = global_mem_read(g_src, g_row - 1, g_col + 1, g_width, g_height);

        s_src[(s_src_row) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row, g_col, g_width, g_height);
        s_src[(s_src_row) * s_src_width + (s_src_col + 1)] = global_mem_read(g_src, g_row, g_col + 1, g_width, g_height);
    } else if (threadIdx.y == 0) {
        // PAD_TOP top rows
        s_src[(s_src_row - 2) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row - 2, g_col, g_width, g_height);
        s_src[(s_src_row - 1) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row - 1, g_col, g_width, g_height);
        s_src[(s_src_row) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row, g_col, g_width, g_height);
    } else if (threadIdx.x == 0) {
        // PAD_LEFT left rows
        s_src[(s_src_row) * s_src_width + (s_src_col - 2)] = global_mem_read(g_src, g_row, g_col - 2, g_width, g_height);
        s_src[(s_src_row) * s_src_width + (s_src_col - 1)] = global_mem_read(g_src, g_row, g_col - 1, g_width, g_height);
        s_src[(s_src_row) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row, g_col, g_width, g_height);
    } else if (threadIdx.y == (blockDim.y - 1)) {
        // PAD_BOTTOM bottom rows
        s_src[(s_src_row) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row, g_col, g_width, g_height);
        s_src[(s_src_row + 1) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row + 1, g_col, g_width, g_height);
    } else if (threadIdx.x == (blockDim.x - 1)) {
        // PAD_RIGHT right rows
        s_src[(s_src_row) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row, g_col, g_width, g_height);
        s_src[(s_src_row) * s_src_width + (s_src_col + 1)] = global_mem_read(g_src, g_row, g_col + 1, g_width, g_height);
    } else {
        // center pixels
        s_src[(s_src_row) * s_src_width + (s_src_col)] = global_mem_read(g_src, g_row, g_col, g_width, g_height);
    }

    // load g_dst into shared memory
    s_dst[(s_dst_row) * s_dst_width + (s_dst_col)] = global_mem_read(g_dst, g_row, g_col, g_width, g_height);

    __syncthreads();

    uint8_t NZ = black_neighbors_around(s_src, s_src_row, s_src_col, s_src_width);
    uint8_t TR_P1 = wb_transitions_around(s_src, s_src_row, s_src_col, s_src_width);
    uint8_t TR_P2 = wb_transitions_around(s_src, s_src_row - 1, s_src_col, s_src_width);
    uint8_t TR_P4 = wb_transitions_around(s_src, s_src_row, s_src_col - 1, s_src_width);
    uint8_t P2 = P2_f(s_src, s_src_row, s_src_col, s_src_width);
    uint8_t P4 = P4_f(s_src, s_src_row, s_src_col, s_src_width);
    uint8_t P6 = P6_f(s_src, s_src_row, s_src_col, s_src_width);
    uint8_t P8 = P8_f(s_src, s_src_row, s_src_col, s_src_width);

    uint8_t thinning_cond_1 = ((2 <= NZ) & (NZ <= 6));
    uint8_t thinning_cond_2 = (TR_P1 == 1);
