 		
 		cudaMemcpy(d_V, V, m*n*sizeof(int), cudaMemcpyHostToDevice);
 
-		dim3 blk(m, m, n);
-		knn<<<1, blk, m*m*sizeof(int)>>>(m, n, k, d_V, d_out, dim_D);
+//		dim3 blk(m, m, n);
+		dim3 blk(m, n);
+		knn<<<m, blk, m*n*sizeof(int)>>>(m, n, k, d_V, d_out, D);
 
 		cudaMemcpy(out, d_out, m*k*sizeof(int), cudaMemcpyDeviceToHost);
 

	int num;
	int dist;
	int is_duplicate;
//	__shared__ int D[m*m];

	computeDist(m, n, V, dim_D);
	__syncthreads();
	// let the first thread select the k-min dist
	i = threadIdx.x;
	if(threadIdx.y == 0)
	{
		for(count=0; count<k; count++)
		{

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

extern __shared__ int block_dim_D[];

__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]);
}

__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;
	//dimention calculation
	computeDimDist(m, n, V);
	__syncthreads();
	//reduce duplications
	base_idx = j*n;
	if(i < j)
	{
		//paralle 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 dim_D needs to be added
		if(n > (n/2)*2)
		{
			dist += block_dim_D[base_idx+n-1];
		}
	}
	D[i*m+j] = dist;
}
