 		const double* points, index_t points_nb,
 		const 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
 		) :
 		vertex_(vertex),
 		vertex_nb_(vertex_nb),

		index_t*		facets_nn, 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){
			g_seeds_information[tid * 4 + 0] = 0.0;
			g_seeds_information[tid * 4 + 1] = 0.0;
			g_seeds_information[tid * 4 + 2] = 0.0;
			g_seeds_information[tid * 4 + 3] = 0.0;
		}

		//load \memory[facet] 3 times.
		int3 facet_index = {
			facets[tid * dim + 0],
			facets[tid * dim + 1],
			facets[tid * dim + 2]
		};

		//load \memory[vertex] 9 times.
		double3 v1 = {
			fetch_double(t_vertex, facet_index.x * dim + 0),
			fetch_double(t_vertex, facet_index.x * dim + 1),
			fetch_double(t_vertex, facet_index.x * dim + 2)
		};
		double3 v2 = {
			fetch_double(t_vertex, facet_index.y * dim + 0),
			fetch_double(t_vertex, facet_index.y * dim + 1),
			fetch_double(t_vertex, facet_index.y * dim + 2)
		};
		double3 v3 = {
			fetch_double(t_vertex, facet_index.z * dim + 0),
			fetch_double(t_vertex, facet_index.z * dim + 1),
			fetch_double(t_vertex, facet_index.z * dim + 2)
		};

		CudaPolygon current_polygon;
		current_polygon.vertex_nb = 3;

		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;

		CudaPolygon current_store = current_polygon;
		//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[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--;

			t_intersection_clip_facet_SR(
				current_polygon,
				current_seed,
				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();

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

	__host__

			cudaMalloc((void**)&dev_points_, DOUBLE_SIZE * points_nb_ * dimension_);
			cudaMalloc((void**)&dev_facets_, sizeof(index_t) * facet_nb_ * dimension_);
			cudaMalloc((void**)&dev_points_nn_, sizeof(index_t) * points_nb_ * k_);
			cudaMalloc((void**)&dev_facets_nn_, sizeof(index_t) * facet_nb_ * f_k);

			//Output result.
			cudaMalloc((void**)&dev_ret_, sizeof(double) * points_nb_ * 4);
