 		double*			points, index_t			points_nb,
 		index_t*		facets, index_t			facets_nb,
 		index_t*		points_nn, index_t			k_p,
-		index_t*		facets_nn, index_t			dim,
+		index_t*		facets_nn, index_t			k_f,
+		index_t			dim,
 		double*			retdata
 		){
-		index_t tid = blockIdx.x * blockDim.x + threadIdx.x;
-		if (tid >= facets_nb) return;
-
-		if (tid >= 0 && tid < points_nb){
+		index_t fid = blockIdx.y * gridDim.x + blockIdx.x;
+		index_t tid = threadIdx.x;
+		if (fid >= facets_nb) return;
+
+		if (fid < points_nb && tid == 0){
 			g_seeds_information[tid * 4 + 0] = 0.0;
 			g_seeds_information[tid * 4 + 1] = 0.0;
 			g_seeds_information[tid * 4 + 2] = 0.0;

		facets_(facets),
		facet_nb_(facets_nb),
		k_(k_p),
		points_nn_(points_nn),
		facets_nn_(facets_nn),
		dimension_(dim),

		}

		//load \memory[facet] 3 times.
		__shared__ int local_facet_index[9];
		if (tid < 9){
			int a = tid / 3, b = tid % 3;
			local_facet_index[tid] = facets[fid * dim + a] * dim + b;
		}
		__syncthreads();
		//load \memory[vertex] 9 times.
		__shared__ double local_vertex[9];

		if (tid < 9){
			local_vertex[tid] = vertex[local_facet_index[tid]];
		}
		__syncthreads();
		
		CudaPolygon current_polygon;
		current_polygon.vertex_nb = 3;

		current_polygon.vertex[0].x = local_vertex[0]; current_polygon.vertex[0].y = local_vertex[1]; current_polygon.vertex[0].z = local_vertex[2]; current_polygon.vertex[0].w = 1.0;
		current_polygon.vertex[1].x = local_vertex[3]; current_polygon.vertex[1].y = local_vertex[4]; current_polygon.vertex[1].z = local_vertex[5]; current_polygon.vertex[1].w = 1.0;
		current_polygon.vertex[2].x = local_vertex[6]; current_polygon.vertex[2].y = local_vertex[7]; current_polygon.vertex[2].z = local_vertex[8]; current_polygon.vertex[2].w = 1.0;

		index_t current_seed = facets_nn[fid * k_f + tid];

		intersection_clip_facet_SR(
			current_polygon,
			current_seed,
			points,
			points_nb,
			points_nn,
			k_p
			);

		//now we get the clipped polygon stored in "polygon", do something.
		/*action(
			current_polygon,
			current_seed
			);*/
		////doesn't have the stack?
		//index_t to_visit[CUDA_Stack_size];
		//index_t to_visit_pos = 0;

		//index_t has_visited[CUDA_Stack_size];
		//index_t has_visited_nb = 0;
		//bool has_visited_flag = false;

		//load \memory[facets_nn] 1 time. 
		/*to_visit[to_visit_pos++] = facets_nn[fid * k_p + tid];
		has_visited[has_visited_nb++] = to_visit[0];*/

		//while (to_visit_pos){
		//	index_t current_seed = to_visit[to_visit_pos - 1];
		//	to_visit_pos--;
		//	
		//	intersection_clip_facet_SR(
		//		current_polygon,
		//		current_seed,
		//		points,
		//		points_nb,
		//		points_nn,
		//		k_p
		//		);

		//	//now we get the clipped polygon stored in "polygon", do something.
		//	action(
		//		current_polygon,
		//		current_seed
		//		);
		//	
		//	//Propagate to adjacent seeds
		//	for (index_t v = 0; v < current_polygon.vertex_nb; ++v)
		//	{
		//		CudaVertex ve = current_polygon.vertex[v];
		//		int ns = ve.neigh_s;
		//		if (ns != -1)
		//		{
		//			for (index_t 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;
		//		}
		//	}	
		//	current_polygon = current_store;
		//}
		__syncthreads();

		/*if (fid == 0 && tid == 0){
			for (index_t i = 0; i < points_nb / 2; ++i){
				retdata[i] = g_seeds_polygon_nb[i];
			}
		}*/
		/*for (index_t i = 0; i < points_nb * 4; ++i){
			retdata[i] = g_seeds_information[i];
		}*/
