 		cudaFree(dev_facets_);
 		cudaFree(dev_points_nn_);
 		cudaFree(dev_facets_nn_);
-		cudaFree(dev_ret);
-
-		if (host_ret != nil){
-			free(host_ret);
-			host_ret = nil;
+		cudaFree(dev_ret_);
+		cudaFree(dev_seeds_info_);
+		cudaFree(dev_seeds_poly_nb);
+		if (host_ret_ != nil){
+			free(host_ret_);
+			host_ret_ = nil;
 		}
 	}
 
 	__host__
 	void CudaRestrictedVoronoiDiagram::copy_back(){
-		cudaMemcpy(host_ret, dev_ret, sizeof(double) * points_nb_ * 4, cudaMemcpyDeviceToHost);
+		cudaMemcpy(host_ret_, dev_ret_, sizeof(double) * points_nb_ * 4, cudaMemcpyDeviceToHost);
 		CheckCUDAError("copy back");
 	}
 
 	__host__
 		void CudaRestrictedVoronoiDiagram::print_return_data(std::string filename) const{
+		for (int i = 0; i < points_nb_; ++i)
+		{
+			if (fabs(host_ret_[i * 4 + 3]) >= 1e-12){
+				host_ret_[i * 4 + 0] /= host_ret_[i * 4 + 3];
+				host_ret_[i * 4 + 1] /= host_ret_[i * 4 + 3];
+				host_ret_[i * 4 + 2] /= host_ret_[i * 4 + 3];
+			}
+		}
 		index_t line_num = 4;
 		std::ofstream f;
 		f.open(filename);
 		for (index_t t = 0; t < points_nb_; ++t){
 			f << std::setprecision(18);
 			f << "point " << t << " ";
-			f << host_ret[t] << " "
-				<< host_ret[t + 1] << " "
-				<< host_ret[t + 2] << " "
-				<< host_ret[t + 3] << " " << std::endl;
+			f << host_ret_[t * 4 + 0] << " "
+				<< host_ret_[t * 4 + 1] << " "
+				<< host_ret_[t * 4 + 2] << " "
+				<< host_ret_[t * 4 + 3] << " " << std::endl;
 		}
 		f.close();
 	}

			dev_facets_, facet_nb_,
			dev_points_nn_, k_,
			dev_facets_nn_, dimension_,
			dev_ret
			);
		CheckCUDAError("kernel function");
		copy_back();

		){
		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 = {
