 					dev_facets_, facet_nb_,
 					dev_points_nn_, k_,
 					dev_facets_nn_, dimension_,
-					dev_ret_, dev_retidx
+					dev_ret_, f_k_
 					);
 				CheckCUDAError("kernel function");
 				copy_back();
-
+				//printf("--------------passsed kernel----------------");
 				//result_print("facetnn.txt", facets_nn_, facet_nb_, 4);
 				//result_print("retdata.txt", host_ret_, points_nb_ * 4, 4);
 				//result_print("retidx.txt", host_retidx, facet_nb_ * 100, 10);
-			
-				for (index_t t = 0; t < points_nb_ * 4; ++t){
-					update_data[t] += host_ret_[t];
-				}
-				
-				for (index_t t = 0; t < facet_nb_; ++t){
-					has_visited[t].insert(facets_nn_[t]);
-					for (index_t s = 0; s < data_size; ++s){
-						int cur_seed = host_retidx[t * data_size + s];
-						if (cur_seed < -1) break;
-						if (cur_seed >= 0 && cur_seed < points_nb_){
-							if (has_visited[t].find(cur_seed) == has_visited[t].end()){
-								has_visited[t].insert(cur_seed);
-								to_visited[t].push(cur_seed);
-							}
-						}
-					}
-				}
-
-				while (!check_task_finished(to_visited)){
-					
-					//put the data into a special array.
-					go_next(next, data_size + 1);
-					cudaMemcpy(dev_next, next, sizeof(int) * facet_nb_ * (data_size + 1), cudaMemcpyHostToDevice);
-					cudaMemset(dev_retidx, -2, sizeof(int) * facet_nb_ * 100);
-					//result_print("next.txt", next, facet_nb_ * (data_size+1), data_size + 1);
-					//go
-					threads = 128;
-					blocks = (facet_nb_ * data_size) / threads + (((facet_nb_ * data_size) % threads) ? 1 : 0);
-					go_rvd << < blocks, threads >> >(
-						dev_vertex_, vertex_nb_,
-						dev_points_, points_nb_,
-						dev_facets_, facet_nb_,
-						dev_points_nn_, k_,
-						dev_next, dimension_, data_size,
-						dev_ret_, dev_retidx
-						);
-
-					cudaMemcpy(host_ret_, dev_ret_, sizeof(double) * points_nb_ * 4, cudaMemcpyDeviceToHost);
-					cudaMemcpy(host_retidx, dev_retidx, sizeof(int) * facet_nb_ * data_size * 10, cudaMemcpyDeviceToHost);
-					CheckCUDAError("copy back");
-					//result_print("go_rvd.txt", host_retidx, facet_nb_ * 100, 10);
-					for (index_t t = 0; t < points_nb_ * 4; ++t){
-						update_data[t] += host_ret_[t];
-					}
-					insert_to_visited(host_retidx, data_size);
-				}
-
-
+		
 				//result_print("retdata.txt", host_ret_, facet_nb_ * 400, 4);
 				//result_print("C:\\Users\\JWhan\\Desktop\\DATA\\retdata.txt", host_ret_, points_nb_ * 4, 4);
 				is_store_ = false;
-				for (int i = 0; i < points_nb_; ++i){
-					double d = update_data[i * 4 + 3];
-					if (fabs(update_data[i * 4 + 3]) > 1e-30){
-						update_data[i * 4 + 0] /= update_data[i * 4 + 3];
-						update_data[i * 4 + 1] /= update_data[i * 4 + 3];
-						update_data[i * 4 + 2] /= update_data[i * 4 + 3];
-
-						x_->set_vertex(&update_data[i * 4], dimension_, i);
-					}
-					else
-					{
-						//std::cout << "point  " << i << " hsa no RVD." << std::endl;
-
-					}
-					//x_->add_vertexd(&host_ret_[i * 4], dimension_);
-				}
-
-				if (is_store_){
-					std::vector<int> sample_facet(points_nb_);
-					std::string name = "C:\\Users\\JWhan\\Desktop\\DATA\\RVD_" + String::to_string(store_filename_counter_) + ".xyz";
-					points_save_xyz(name, *x_, sample_facet);
-					store_filename_counter_++;
-				}
-				//update_points();
+				
+				update_points();
 				iter_watcher.stop();
 				iter_watcher.synchronize();
 				iter_watcher.print_elaspsed_time(std::cout);

	}

	void CudaRestrictedVoronoiDiagram::store_f_neighrbors_CB(index_t v){
		index_t nb = 1;

		// Allocated on the stack(more thread-friendly and 
		// no need to deallocate)

		current_polygon.vertex[2].y = vertex[facet_index.z * dim + 1];
		current_polygon.vertex[2].z = vertex[facet_index.z * dim + 2];
		current_polygon.vertex[2].w = 1.0;
		
		index_t current_seed = facets_nn[tid];
		
		intersection_clip_facet_SR(
			current_polygon,
			current_seed,
