kernels.cu

#include "cuPrintf.cu"
#include <stdio.h>

__device__ __host__ int distance(int x1, int y1, int x2, int y2) 
{
	return (x1 - x2) * (x1 - x2) + (y1 - y2) * (y1 - y2); 
}

/* The CPU version of Voronoi Diagram computation */ 
void cpuVoronoiDiagram(Point2D *points, int *output, int noPoints, int width, int height) 
{
	int id = 0; 
	for (int j = 0; j < height; j++) 
		for (int i = 0; i < width; i++) {
			int best = 0; 
			int minDist = distance(i, j, points[0].x, points[0].y);

			for (int k = 1; k < noPoints; k++) {
				int myDist = distance(i, j, points[k].x, points[k].y); 

				if (myDist < minDist) {
					minDist = myDist; 
					best = k;
				}
			}

			output[id++] = best; 
		}
}

/* GPU Version of Voronoi diagram computation 
 * Naive implementation
 */
__global__ void kernelNaiveVoronoiDiagram(Point2D *points, int *result, int noPoints, int width) 
{
	// 1. Get the coordinate of this thread
	int iy = blockDim.y*blockIdx.y+threadIdx.y;
	int ix = blockDim.x*blockIdx.x+threadIdx.x;
	int idx = iy * width + ix;

	// 2. Initialize

	// 3. Find the nearest site
	int best = 0; 
	int minDist = distance(ix, iy, points[0].x, points[0].y);

	for (int k = 1; k < noPoints; k++) {
		int myDist = distance(ix, iy, points[k].x, points[k].y); 

		if (myDist < minDist) {
			minDist = myDist; 
			best = k;
		}
	}

	// 4. Output the result
	result[idx] = best; 
}

/* GPU Version of Voronoi diagram computation
 * Using shared memory
 * We assume for simplicity that noPoints is divisible by TILE
 */
__global__ void kernelSharedVoronoiDiagram(Point2D *points, int *result, int noPoints, int width) 
{
	// 1. Get the coordinate of this thread
	int iy = blockDim.y*blockIdx.y+threadIdx.y;
	int ix = blockDim.x*blockIdx.x+threadIdx.x;
	int idx = iy * width + ix;
	int tdx = threadIdx.x * 16 + threadIdx.y; 

	// 2. Initialize
	__shared__ int index_x[256];
	__shared__ int index_y[256];
	__syncthreads();
	
	// 3. Find the nearest site
		// 3.1 Each thread read one point into the shared array
		index_x[tdx]=points[tdx].x;
		index_y[tdx]=points[tdx].y;

		// 3.2 Make sure everyone has finished reading
		__syncthreads();

		// 3.3 Check with all points in the shared array
		int best = 0; 
		int minDist = distance(ix, iy, index_x[0], index_y[0]);
		for (int k = 1; k < noPoints; k++) {
			int myDist = distance(ix, iy, index_x[k],index_y[k]); 

			if (myDist < minDist) {
				minDist = myDist; 
				best = k;
			}
		}
		// 3.4 Make sure everyone has finished using the shared array
	// 4. Output the result
        result[idx]=best;
}