 
     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)];
-    uint8_t* s_equ = &s_data[((blockDim.x + PAD_LEFT + PAD_RIGHT) * (blockDim.y + PAD_TOP + PAD_BOTTOM)) + (blockDim.x * blockDim.y)];
 
     int g_row = blockIdx.y * blockDim.y + threadIdx.y;
     int g_col = blockIdx.x * blockDim.x + threadIdx.x;

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

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

    // load g_src & g_dst 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);

    do {
        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_dst_size = block_dim.x * block_dim.y * sizeof(uint8_t);
        int skeletonize_pass_shared_mem_size = skeletonize_pass_s_src_size + skeletonize_pass_s_dst_size;
        skeletonize_pass<<<grid_dim, block_dim, skeletonize_pass_shared_mem_size>>>(g_src_data, g_dst_data, g_equ_data, (*src_bitmap)->width, (*src_bitmap)->height);
        gpuErrchk(cudaPeekAtLastError());
        gpuErrchk(cudaDeviceSynchronize());
