Brute Force Exact Euclidian Distance Transform in CUDA

Hi!

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

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

  float minv = INFINITY;

  if (img[i] > 0)
  {
    minv = 0.0f;
  }
  else
  {
    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;
}

Performance

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 https://github.com/victormatheus/DT-GPU

About these ads