Brute Force Exact Euclidian Distance Transform in CUDA


Following a discussion in Reddit about the Distance Transform in GPU, I decided to post a implementation I made some time ago.

It’s the brute force euclidian distance transform. Basically, in a binary image, for each pixel in the foreground we verify what is the 2D euclidian distance to the nearest pixel in the background.

Here is the source of the kernel:

#define BLOCK_SIZE 256

__global__ void euclidian_distance_transform_kernel(
  const unsigned char* img, float* dist, int w, int h)
  const int i = blockIdx.x*blockDim.x + threadIdx.x;
  const int N = w*h;

  if (i >= N)

  int cx = i % w;
  int cy = i / w;

  float minv = INFINITY;

  if (img[i] > 0)
    minv = 0.0f;
    for (int j = 0; j < N; j++)
        if (img[j] > 0)
          int x = j % w;
          int y = j / w;
          float d = sqrtf( powf(float(x-cx), 2.0f) + powf(float(y-cy), 2.0f) );
          if (d < minv) minv = d;

  dist[i] = minv;


35.8 seconds for a 1 megapixel image in a Tesla C2050 GPU. Terrible result.

I suppose a good hardware can’t save a bad algorithm after all 😉

The code is at

This entry was posted in gpu, image processing. Bookmark the permalink.

Leave a Reply

Fill in your details below or click an icon to log in: Logo

You are commenting using your account. Log Out / Change )

Twitter picture

You are commenting using your Twitter account. Log Out / Change )

Facebook photo

You are commenting using your Facebook account. Log Out / Change )

Google+ photo

You are commenting using your Google+ account. Log Out / Change )

Connecting to %s