 	int dist;
 	int is_duplicate;
 
-	computeDist(m, n, V, D);
-	__syncthreads();
 	// find the k nearest neighbors of the point with index = blockIdx.x
 	i = blockIdx.x;
 	// let the first thread select the k-min distance

#define INIT_MAX 100
void showResult(int m, int k, int *out);

extern __shared__ int block_dim_D[];

// compute the square of distance per dimension
// the kth dimension of the ith point and jth point
__device__ void computeDimDist(int m, int n, int *V)
{
	int i = blockIdx.x;
   	int j = threadIdx.x;
	int k = threadIdx.y;
	block_dim_D[j*n+k] = (V[i*n+k]-V[j*n+k])*(V[i*n+k]-V[j*n+k]);
}

// compute the square of distance of the ith point and jth point
__device__ void computeDist(int m, int n, int *V, int *D)
{
	int i = blockIdx.x;
   	int j = threadIdx.x;
	int k = threadIdx.y;
	int s;
	int base_idx;
	int dist = 0;
	// calculate the square of distance per dimensions
	computeDimDist(m, n, V);
	__syncthreads();
	base_idx = j*n;
	// reduce duplicated calculations since d(i, j) = d(j, i)
	// also, we do not consider the trivial case of d(i, i) = 0
	// so we only compute the square distance when i < j 
	if(i < j)
	{
		// use parallel reduction
		for(s=n/2; s>0; s>>=1)
		{
			if(k < s)
			{
				block_dim_D[base_idx+k] += block_dim_D[base_idx+k+s];
			}
		}
		__syncthreads();
		dist = block_dim_D[base_idx];
		// when n is odd, the last element of block_dim_D needs to be added
		if(n > (n/2)*2)
		{
			dist += block_dim_D[base_idx+n-1];
		}
	}
	D[i*m+j] = dist;
}

// compute the k nearest neighbors
__global__ void knn(int m, int n, int k, int *V, int *out, int *D)
{
	int i,j;
	int temp;

		// record event
		cudaEventRecord(start);
		// launch knn() kernel on GPU
		computeDist<<<grid, n, n*sizeof(int)>>>(m, n, d_V, D);
		cudaDeviceSynchronize();
		knn<<<m, m, m*sizeof(int)>>>(m, n, k, d_V, D, d_out);
		// record event and synchronize
		cudaEventRecord(stop);
		cudaEventSynchronize(stop);
