 
 	CheckCUDAError("pass data back to CPU");
 
-	/*for (int i = 0; i < points_nb; ++i)
+	for (int i = 0; i < points_nb; ++i)
 	{
 		if (host_seedsInfo[i * 4 + 3] != 0){
 			host_seedsInfo[i * 4 + 0] /= host_seedsInfo[i * 4 + 3];
 			host_seedsInfo[i * 4 + 1] /= host_seedsInfo[i * 4 + 3];
 			host_seedsInfo[i * 4 + 2] /= host_seedsInfo[i * 4 + 3];
 		}
-	}*/
-
+	}
+	std::ofstream fileout("1.txt");
 	for (int i = 0; i < points_nb; ++i)
 	{
-		printf("Line %d : x : %.17lf, y : %.17lf, z : %.17lf, w : %.17lf\n", i, host_seedsInfo[i * 4 + 0], host_seedsInfo[i * 4 + 1], host_seedsInfo[i * 4 + 2], host_seedsInfo[i * 4 + 3]);
+		//printf("Line %d : x : %.17lf, y : %.17lf, z : %.17lf, w : %.17lf\n", i, host_seedsInfo[i * 4 + 0], host_seedsInfo[i * 4 + 1], host_seedsInfo[i * 4 + 2], host_seedsInfo[i * 4 + 3]);
+		fileout << "Line " << i << ':' << setprecision(16) << host_seedsInfo[i * 4 + 0] << ' ' << host_seedsInfo[i * 4 + 1] << ' ' << host_seedsInfo[i * 4 + 2] << ' ' << host_seedsInfo[i * 4 + 3] << endl;
 	}
 
 	free(host_seedsInfo);

			MyAtomicAdd(&SeedsInformation[current_seed * 4 + 2], temp_pos.z);
			MyAtomicAdd(&SeedsInformation[current_seed * 4 + 3], current_weight);
		}

		if (tid == 0)
		{
			for (int m = 0; m < seeds_nb * 4; ++m)
				ret_seeds[m] = m;

			ret_seeds[0] = v1.x;
			ret_seeds[1] = v1.y;
			ret_seeds[2] = v1.z;

			for (int m = 0; m < vertex_nb * 4; ++m)
			{
				ret_seeds[4 + m] = polygon[m];
			}
			ret_seeds[21] = vertex_nb;
		}
		if (i == 1)
		return;
	}
	__syncthreads();

	
	
	for (int i = 0; i < seeds_nb * 4; ++i)
	{
		ret_seeds[i] = SeedsInformation[i];
	}
	
	free(polygon);
	return;
}


/*
*/
__device__
void intersection_clip_facet_with_knn(double* polygon_pointer, int& vertex_nb, int current_seed, double* seeds_pointer, int seeds_nb, int *seeds_neighbor_index, int k
, int tid, int pass)
{
	
	//set a buffer pointer to store the polygon
	double polygon_buffer[40];
	int polygon_buffer_nb = 0;
	
	for (int i = 0; i < k; ++i)
	{
		int j = seeds_neighbor_index[current_seed * k + i];
