 
 	int threads = 512;
 	int blocks = mesh_facet_nb / threads + ((mesh_facet_nb % threads) ? 1 : 0);
-
+	
+	cudaEvent_t start, stop;
+	cudaEventCreate(&start);
+	cudaEventCreate(&stop);
+	float elapsed_time;
+	cudaEventRecord(start, 0);
 	compute_RVD_with_knn << <threads, blocks >> >(dev_seeds_pointer, points_nb, dev_mesh_vertex_pointer, mesh_vertex_nb, dev_mesh_facet_index, mesh_facet_nb, dev_facet_center_neighbors,
 		dev_seeds_neighbors, 10, ret_seeds, test_seeds);
 

ping store the result;
*/
__device__
void swap_polygons(double* ping, int& ping_nb, double* pong, int& pong_nb)
{
	//!!! can be accerate
	//ȷǷʹmemset memcpy
	double temp;
	int temp_nb = ping_nb;

	ping_nb = pong_nb;
	pong_nb = temp_nb;

	for (int i = 0; i < 40; ++i)
	{
		temp = ping[i];
		ping[i] = pong[i];
		pong[i] = temp;
	}
}

/*
*/
__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];
		if (current_seed != j)
		{
			clip_by_plane(polygon_pointer, vertex_nb, polygon_buffer, polygon_buffer_nb, current_seed, j, seeds_pointer, seeds_nb);
			swap_polygons(polygon_pointer, vertex_nb, polygon_buffer, polygon_buffer_nb);
		}
	}
	return;
}

/*
	device function
	compute the intersection of a facet and a seed(a cell)

	input:

*/
__device__
void intersection_clip_facet(double* polygon_pointer, int& vertex_nb, int current_seed, double* seeds_pointer, int seeds_nb)
{
	double3 current_seed_pos = { seeds_pointer[current_seed * 3 + 0],
		seeds_pointer[current_seed * 3 + 1],
		seeds_pointer[current_seed * 3 + 2] };

	int* seeds_neighbors = (int*)malloc(sizeof(int) * 20);
	findNearestPoints(current_seed_pos, seeds_pointer, seeds_nb, seeds_neighbors, 20);

	////set a buffer pointer to store the polygon
	double* polygon_buffer = (double*)malloc(sizeof(double) * 10 * 4);
	int polygon_buffer_nb = 0;

	for (int i = 0; i < 20; ++i)
	{
		int j = seeds_neighbors[i];

		if (current_seed != j)
		{ 
			clip_by_plane(polygon_pointer, vertex_nb, polygon_buffer, polygon_buffer_nb, current_seed, j, seeds_pointer, seeds_nb);
			swap_polygons(polygon_pointer, vertex_nb, polygon_buffer, polygon_buffer_nb);
		}
	}
	free(seeds_neighbors);
	free(polygon_buffer);
}

__global__

		mesh_vertex[facet_index.z * 3 + 1],
		mesh_vertex[facet_index.z * 3 + 2] };

	Cuda_Polygon current_polygon;
	current_polygon.vertex_nb = 3;

	//intialize the polygon with the 3 vertex of the current facet
	/*
	polygon pointer can be apart by several vertex
	a vertex is made up with x,y,z,w
	w : the weight of a vertex
	*/
	
	current_polygon.vertex[0].x = v1.x; current_polygon.vertex[0].y = v1.y; current_polygon.vertex[0].z = v1.z; current_polygon.vertex[0].w = 1.0;
	current_polygon.vertex[1].x = v2.x; current_polygon.vertex[1].y = v2.y; current_polygon.vertex[1].z = v2.z; current_polygon.vertex[1].w = 1.0;
	current_polygon.vertex[2].x = v3.x; current_polygon.vertex[2].y = v3.y; current_polygon.vertex[2].z = v3.z; current_polygon.vertex[2].w = 1.0;
	
	/*if (tid == 0)
	{
		ret_seeds[0] = current_polygon.vertex[0].x;
		ret_seeds[1] = current_polygon.vertex[0].y;
		ret_seeds[2] = current_polygon.vertex[0].z;
		ret_seeds[3] = current_polygon.vertex[0].w;
		ret_seeds[4] = current_polygon.vertex[0].neigh_s;
	}*/
	
	//doesn't have the stack?
	int to_visit[CUDA_Stack_size];
	int to_visit_pos = 0;

	int has_visited[CUDA_Stack_size];
	int has_visited_nb = 0;
	bool has_visited_flag = false;
	
	to_visit[to_visit_pos++] = facet_center_neighbor_index[tid];
	has_visited[has_visited_nb++] = to_visit[0];
	while (to_visit_pos){
		int current_seed = to_visit[to_visit_pos - 1];
		to_visit_pos--;

		intersection_clip_facet_with_knn(current_polygon, current_seed, seeds_pointer, seeds_nb, seeds_neighbor_index, k);
		
		//use the RVD
		//action(current_polygon, current_seed);

		//now we get the clipped polygon stored in "polygon"
		//take care of the sychonize
		//change the polygon data into "weight" and "position"
		
		//Propagate to adjacent seeds
		for (int v = 0; v < current_polygon.vertex_nb; ++v)
		{
			Cuda_Vertex ve = current_polygon.vertex[v];
			int ns = ve.neigh_s;
			if (ns != -1)
			{
				for (int ii = 0; ii < has_visited_nb; ++ii)
				{
					//if the neighbor seed has clipped the polygon
					//the flag should be set "true"
					if (has_visited[ii] == ns)
						has_visited_flag = true;
				}
				//the neighbor seed is new!
				if (!has_visited_flag)
				{
					to_visit[to_visit_pos++] = ns;
					has_visited[has_visited_nb++] = ns;
				}
				has_visited_flag = false;
			}
		}
	}

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