 		float time;
 		// get event elapsed time
 		cudaEventElapsedTime(&time, start, stop);
+
+		// copy result back to host
+		cudaMemcpy(out, d_out, m*k*sizeof(int), cudaMemcpyDeviceToHost);
+
+//		// record event and synchronize
+//		cudaEventRecord(stop);
+//		cudaEventSynchronize(stop);
+//		float time;
+//		// get event elapsed time
+//		cudaEventElapsedTime(&time, start, stop);
 
 		showResult(m, k, out);
 		if(m == 1024) {

void showD(int m, int *D)
{
	int i,j;
printf("D:\n");
	for(i=0; i<m; i++)
	{
		for(j=0; j<m; j++)

{
	int i = blockIdx.x;
  	int tid = threadIdx.x;

	int s = blockDim.x/2;
	int resultValue = INIT_MAX;
	int resultIndex = INIT_MAX;
	int indexBase = (m<MAX_PTRNUM_IN_SMEM)? m: MAX_PTRNUM_IN_SMEM;
	
	for(int num=0; num<m; num+=MAX_PTRNUM_IN_SMEM)
	{
		for(int j=tid; j<indexBase; j+=blockDim.x)
		{
			if(j+num == i)
			{
				SMem[j] = INIT_MAX;
			}
			else
			{
				SMem[j] = D[i*m+num+j];
			}
			//index
			SMem[indexBase+j] = j+num;
			__syncthreads();
/*
			if(tid < count)
			{
				if(out[i*k+tid]-num>0 && out[i*k+tid]-num < indexBase)
				{
					SMem[ out[i*k+tid]-num ] = INIT_MAX;
				}
			}
			__syncthreads();
*/
		}
//		__syncthreads();
		for(int j=0; j<count; j++)
		{
			if(out[i*k+j]-num>=0 && out[i*k+j]-num<indexBase)
			{
				SMem[ out[i*k+j]-num ] = INIT_MAX;
			}
			__syncthreads();
		}
		__syncthreads();
/*
		if(tid < count)
		{
			if(out[i*k+tid]-num>0 && out[i*k+tid]-num < indexBase)
			{
				SMem[ out[i*k+tid]-num ] = INIT_MAX;
			}
//			SMem[ out[i*k+tid] ] = INIT_MAX;
		}
		__syncthreads();
*/
		for(s=indexBase/2; s>0; s>>=1) 
		{
			for(int j=tid; j<indexBase; j+=blockDim.x)
			{
				if(j < s) 
				{
					if(SMem[j] == SMem[j+s])
					{
						if(SMem[indexBase+j] > SMem[indexBase+j+s])
						{
							SMem[indexBase+j] = SMem[indexBase+j+s];
						}
					}
					else if(SMem[j] > SMem[j+s])
					{
						SMem[j] = SMem[j+s];
						SMem[indexBase+j] = SMem[indexBase+j+s];
					}
				}
				__syncthreads();
			}
		}
	
		if(resultValue == SMem[0])
		{
			if(resultIndex > SMem[indexBase])
			{
				resultIndex = SMem[indexBase];
			}
		} 
		else if (resultValue > SMem[0])
		{
			resultValue = SMem[0];
			resultIndex = SMem[indexBase];
		}
		__syncthreads();
	}
	return resultIndex;

}

// compute the k nearest neighbors
