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

	{
		for(j=0; j<k; j++)
		{
				printf("%d", out[i*k+j]);
				if(j == k-1)
				{
						printf("\n");
				}
				else
				{
						printf(" ");
				}
		}
	} 
} 


__device__ void computeDist(int m, int n, int *V)
{
	int i=threadIdx.x;
   	int j=threadIdx.y;
	int k;
	int dist = 0;
	//reduce duplications
	if(i < j)
	{
		for(k=0; k<n; k++)
		{
			dist += (V[i*n+k]-V[j*n+k])*(V[i*n+k]-V[j*n+k]);
		}
	}
	D[i*m+j] = dist;
}

//paralle reduction?
//cannot use D now!
__device__ int prmin(int m, int *D, int *R)
{
	int j = threadIdx.y;
	int s = blockDim.y/2;
	R[j] = j;
	__syncthreads();
	for(s=blockDim.y; s>0; s>>=1)
	{
		if(j < s)
		{
			D[j] = D[j]<D[j+s]? D[j]: D[j+s]; 
			R[j] = R[j]<R[j+s]? R[j]: R[j+s]; 
		}
		__syncthreads();
	}
	return R[0];
}

__global__ void knn(int m, int n, int k, int *V, int *out)
