 		// 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, R);
+		//test
+		cudaMemcpy(h_D, D, m*m*sizeof(int), cudaMemcpyDeviceToHost);
+		showD(m, h_D);
+		knn<<<m, m, m*sizeof(int)>>>(m, n, k, d_V, D, d_out);
 		// record event and synchronize
 		cudaEventRecord(stop);
 		cudaEventSynchronize(stop);

	}
}

__device__ int findMin(int m, int n, int k, int count, int *D, int *out, int *R)
{
//	__shared__ int R[1000];
	int i = blockIdx.x;
  	int j = threadIdx.x;
	int s = blockDim.x/2;
	int num;
	int baseIdx = i*m;
	initSM(m, D);
	__syncthreads();
	R[baseIdx+j] = j;
	__syncthreads();
	SM[i] = INIT_MAX;
	__syncthreads();

		SM[ out[i*k+num] ] = INIT_MAX;
	}
	__syncthreads();
	for(s=m/2; s>0; s>>=1) 
	{
		// check whether the jth point is the same point as the ith one
		// or has already in the k-nn list
		if(j < s) 
		{
			if(SM[j] == SM[j+s])
			{
				if(R[j] > R[j+s])
				{
					R[j] = R[j+s];
				}
			}
			else if(SM[j] > SM[j+s])
			{
				SM[j] = SM[j+s];
				R[j] = R[j+s];
			}
			__syncthreads();
		}
	}
	// when m is odd, the last element of SM needs to be compared
	if(m > (m/2)*2)
	{
		if(SM[0] > SM[m-1])
		{
			R[0] = R[m-1];
		}
	}
	__syncthreads();
	return R[0];
}

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