 // 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_src[];
+    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_row = threadIdx.y + PAD_TOP;
-    int s_col = threadIdx.x + PAD_LEFT;
-    int s_width = blockDim.x + PAD_LEFT + PAD_RIGHT;
-
-    // load data into shared memory
+    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_row - 2) * s_width + (s_col - 2)] = global_mem_read(g_src, g_row - 2, g_col - 2, g_width, g_height);
-        s_src[(s_row - 2) * s_width + (s_col - 1)] = global_mem_read(g_src, g_row - 2, g_col - 1, g_width, g_height);
-        s_src[(s_row - 2) * s_width + (s_col)] = global_mem_read(g_src, g_row - 2, g_col, g_width, g_height);
-
-        s_src[(s_row - 1) * s_width + (s_col - 2)] = global_mem_read(g_src, g_row - 1, g_col - 2, g_width, g_height);
-        s_src[(s_row - 1) * s_width + (s_col - 1)] = global_mem_read(g_src, g_row - 1, g_col - 1, g_width, g_height);
-        s_src[(s_row - 1) * s_width + (s_col)] = global_mem_read(g_src, g_row - 1, g_col, g_width, g_height);
-
-        s_src[(s_row) * s_width + (s_col - 2)] = global_mem_read(g_src, g_row, g_col - 2, g_width, g_height);
-        s_src[(s_row) * s_width + (s_col - 1)] = global_mem_read(g_src, g_row, g_col - 1, g_width, g_height);
-        s_src[(s_row) * s_width + (s_col)] = global_mem_read(g_src, g_row, g_col, g_width, g_height);
+        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_row) * s_width + (s_col - 2)] = global_mem_read(g_src, g_row, g_col - 2, g_width, g_height);
-        s_src[(s_row) * s_width + (s_col - 1)] = global_mem_read(g_src, g_row, g_col - 1, g_width, g_height);
-        s_src[(s_row) * s_width + (s_col)] = global_mem_read(g_src, g_row, g_col, g_width, g_height);
-
-        s_src[(s_row + 1) * s_width + (s_col - 2)] = global_mem_read(g_src, g_row + 1, g_col - 2, g_width, g_height);
-        s_src[(s_row + 1) * s_width + (s_col - 1)] = global_mem_read(g_src, g_row + 1, g_col - 1, g_width, g_height);
-        s_src[(s_row + 1) * s_width + (s_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);
+
+        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_row) * s_width + (s_col)] = global_mem_read(g_src, g_row, g_col, g_width, g_height);
-        s_src[(s_row) * s_width + (s_col + 1)] = global_mem_read(g_src, g_row, g_col + 1, g_width, g_height);
-
-        s_src[(s_row + 1) * s_width + (s_col)] = global_mem_read(g_src, g_row + 1, g_col, g_width, g_height);
-        s_src[(s_row + 1) * s_width + (s_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);
+
+        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_row - 2) * s_width + (s_col)] = global_mem_read(g_src, g_row - 2, g_col, g_width, g_height);
-        s_src[(s_row - 2) * s_width + (s_col + 1)] = global_mem_read(g_src, g_row - 2, g_col + 1, g_width, g_height);
-
-        s_src[(s_row - 1) * s_width + (s_col)] = global_mem_read(g_src, g_row - 1, g_col, g_width, g_height);
-        s_src[(s_row - 1) * s_width + (s_col + 1)] = global_mem_read(g_src, g_row - 1, g_col + 1, g_width, g_height);
-
-        s_src[(s_row) * s_width + (s_col)] = global_mem_read(g_src, g_row, g_col, g_width, g_height);
-        s_src[(s_row) * s_width + (s_col + 1)] = global_mem_read(g_src, g_row, 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 - 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_row - 2) * s_width + (s_col)] = global_mem_read(g_src, g_row - 2, g_col, g_width, g_height);
-        s_src[(s_row - 1) * s_width + (s_col)] = global_mem_read(g_src, g_row - 1, g_col, g_width, g_height);
-        s_src[(s_row) * s_width + (s_col)] = global_mem_read(g_src, g_row, g_col, 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)] = 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_row) * s_width + (s_col - 2)] = global_mem_read(g_src, g_row, g_col - 2, g_width, g_height);
-        s_src[(s_row) * s_width + (s_col - 1)] = global_mem_read(g_src, g_row, g_col - 1, g_width, g_height);
-        s_src[(s_row) * s_width + (s_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 - 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_row) * s_width + (s_col)] = global_mem_read(g_src, g_row, g_col, g_width, g_height);
-        s_src[(s_row + 1) * s_width + (s_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);
+        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_row) * s_width + (s_col)] = global_mem_read(g_src, g_row, g_col, g_width, g_height);
-        s_src[(s_row) * s_width + (s_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) * 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_row) * s_width + (s_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)] = 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_row, s_col, s_width);
-    uint8_t TR_P1 = wb_transitions_around(s_src, s_row, s_col, s_width);
-    uint8_t TR_P2 = wb_transitions_around(s_src, s_row - 1, s_col, s_width);
-    uint8_t TR_P4 = wb_transitions_around(s_src, s_row, s_col - 1, s_width);
-    uint8_t P2 = P2_f(s_src, s_row, s_col, s_width);
-    uint8_t P4 = P4_f(s_src, s_row, s_col, s_width);
-    uint8_t P6 = P6_f(s_src, s_row, s_col, s_width);
-    uint8_t P8 = P8_f(s_src, s_row, s_col, s_width);
+    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);

    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);
}


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