     pad_binary_bitmap(&dst_bitmap, BINARY_WHITE, padding);
 
     // allocate memory on device
-    // TODO
+    uint8_t* d_src_data = NULL;
+    uint8_t* d_dst_data = NULL;
+    unsigned int data_size = src_bitmap->width * src_bitmap->height * sizeof(uint8_t);
+    cudaError d_src_malloc_success = cudaMalloc((void**) &d_src_data, data_size);
+    cudaError d_dst_malloc_success = cudaMalloc((void**) &d_dst_data, data_size);
+    assert((d_src_malloc_success == cudaSuccess) && "Could not allocate memory for d_src_data");
+    assert((d_dst_malloc_success == cudaSuccess) && "Could not allocate memory for d_dst_data");
+
+    // send data to device
+    cudaMemcpy(d_src_data, src_bitmap->data, data_size, cudaMemcpyHostToDevice);
+    cudaMemcpy(d_dst_data, dst_bitmap->data, data_size, cudaMemcpyHostToDevice);
 
     // iterative thinning algorithm
     unsigned int iterations = 0;
-    // do {
-    //     skeletonize_pass<<<grid_dim, block_dim>>>(src_bitmap->data, dst_bitmap->data, src_bitmap->width, src_bitmap->height, padding);
-    //     swap_bitmaps(&src_bitmap, &dst_bitmap);
+    do {
+        skeletonize_pass<<<grid_dim, block_dim>>>(d_src_data, d_dst_data, src_bitmap->width, src_bitmap->height, padding);
 
-    //     iterations++;
-    //     printf(".");
-    //     fflush(stdout);
-    // } while (!are_identical_bitmaps(src_bitmap, dst_bitmap));
+        // bring data back from device
+        cudaMemcpy(src_bitmap->data, d_src_data, data_size, cudaMemcpyDeviceToHost);
+        cudaMemcpy(dst_bitmap->data, d_dst_data, data_size, cudaMemcpyDeviceToHost);
+
+        swap_bitmaps((void**) &d_src_data, (void**) &d_dst_data);
+
+        iterations++;
+        printf(".");
+        fflush(stdout);
+    } while (!are_identical_bitmaps(src_bitmap, dst_bitmap));
+
+    cudaMemcpy(dst_bitmap->data, d_dst_data, data_size, cudaMemcpyDeviceToHost);
 
     // Remove extra padding that was added to the images (don't care about
     // src_bitmap, so only need to unpad dst_bitmap)

    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;

    dst[row * width + col] = BINARY_WHITE + ((1 - thinning_cond_ok) * src[row * width + col]);
}

// Computes the number of white to black transitions around a pixel.
__device__ uint8_t wb_transitions_around(uint8_t* data, unsigned int row, unsigned int col, unsigned int width) {
    uint8_t count = 0;

    count += ( (P2(data, row, col, width) == BINARY_WHITE) & (P3(data, row, col, width) == BINARY_BLACK) );
    count += ( (P3(data, row, col, width) == BINARY_WHITE) & (P4(data, row, col, width) == BINARY_BLACK) );
    count += ( (P4(data, row, col, width) == BINARY_WHITE) & (P5(data, row, col, width) == BINARY_BLACK) );
    count += ( (P5(data, row, col, width) == BINARY_WHITE) & (P6(data, row, col, width) == BINARY_BLACK) );
    count += ( (P6(data, row, col, width) == BINARY_WHITE) & (P7(data, row, col, width) == BINARY_BLACK) );
    count += ( (P7(data, row, col, width) == BINARY_WHITE) & (P8(data, row, col, width) == BINARY_BLACK) );
    count += ( (P8(data, row, col, width) == BINARY_WHITE) & (P9(data, row, col, width) == BINARY_BLACK) );
    count += ( (P9(data, row, col, width) == BINARY_WHITE) & (P2(data, row, col, width) == BINARY_BLACK) );

    return count;
}

#define PAD_BOTTOM 1
#define PAD_RIGHT 1

#define P2(d_data, row, col, width) ((d_data)[((row)-1) * width +  (col)   ])
#define P3(d_data, row, col, width) ((d_data)[((row)-1) * width + ((col)-1)])
#define P4(d_data, row, col, width) ((d_data)[ (row)    * width + ((col)-1)])
#define P5(d_data, row, col, width) ((d_data)[((row)+1) * width + ((col)-1)])
#define P6(d_data, row, col, width) ((d_data)[((row)+1) * width +  (col)   ])
#define P7(d_data, row, col, width) ((d_data)[((row)+1) * width + ((col)+1)])
#define P8(d_data, row, col, width) ((d_data)[ (row)    * width + ((col)+1)])
#define P9(d_data, row, col, width) ((d_data)[((row)-1) * width + ((col)+1)])

// Computes the number of black neighbors around a pixel.
__device__ uint8_t black_neighbors_around(uint8_t* d_data, unsigned int row, unsigned int col, unsigned int width) {
    uint8_t count = 0;

    count += (P2(d_data, row, col, width) == BINARY_BLACK);
    count += (P3(d_data, row, col, width) == BINARY_BLACK);
    count += (P4(d_data, row, col, width) == BINARY_BLACK);
    count += (P5(d_data, row, col, width) == BINARY_BLACK);
    count += (P6(d_data, row, col, width) == BINARY_BLACK);
    count += (P7(d_data, row, col, width) == BINARY_BLACK);
    count += (P8(d_data, row, col, width) == BINARY_BLACK);
    count += (P9(d_data, row, col, width) == BINARY_BLACK);

    return count;
}
