 
 		// Phase 3
 		// Collassa gli alberi union-find sulle radici
-		PathCompression << <grid_size_, block_size_ >> > (d_img_, d_labels_, d_img_labels_);
+		PathCompression << <grid_size_, block_size_ >> > (d_img_, d_img_labels_);
+		ZeroBackground<<<grid_size_, block_size_ >> > (d_img_, d_img_labels_);
 
 		cudaDeviceSynchronize();
-
-		// Memory transfer Device -> Host
-		// d_img_labels_.download(img_labels_);
-
-		// GPU dealloc
-		// d_img_.release();
-		d_labels_.release();
-		// d_img_labels_.release();
 	}
 
 
 private:
 	// Provo a contare solamente l'allocazione delle strutture dati in Gpu ($)
 	double Alloc() {
-		// Memory allocation for the output image
-		//perf_.start();
-		// img_labels_ = cv::Mat1i(img_.size());
-		//memset(img_labels_.data, 0, img_labels_.dataend - img_labels_.datastart);
-		//perf_.stop();
-		//double t = perf_.last();
-		//perf_.start();
-		//memset(img_labels_.data, 0, img_labels_.dataend - img_labels_.datastart);
-		//perf_.stop();
-		//double ma_t = t - perf_.last();
-		//perf_.start();
-		// Memory allocation for GPU data structures
-		// d_img_.create(img_.size(), CV_8UC1);
-		/*$*/ perf_.start();
-		d_labels_.create(img_.size(), CV_32SC1);
-		d_img_labels_.create(img_.size(), CV_32SC1);
-		perf_.stop();
-		// Initialization of class variables
-		// Metti dei define al posto di questi
-		grid_size_.x = (img_.cols + BLOCK_COLS - 1) / BLOCK_COLS;
-		grid_size_.y = (img_.rows + BLOCK_ROWS - 1) / BLOCK_ROWS;
-		grid_size_.z = 1;
-		block_size_.x = BLOCK_COLS;
-		block_size_.y = BLOCK_ROWS;
-		block_size_.z = 1;
-		// perf_.stop();
-		// return perf_.last() + ma_t;
-		/*$*/ return perf_.last();
-	}
-
-	// $
+		perf_.start();
+		d_img_labels_.create(d_img_.size(), CV_32SC1);
+		perf_.stop();
+		return perf_.last();
+	}
+
 	double Dealloc() {
-		// d_img_.release();
-		perf_.start();
-		d_labels_.release();
-		perf_.stop();
-		// d_img_labels_.release();
+		perf_.start();
+		perf_.stop();
 		return perf_.last();
 	}
 

		perf_.stop();
		perf_.store(Step(StepType::SECOND_SCAN), perf_.last());

		//perf_.start();
		//MemoryTransferDeviceToHost();
		//perf_.stop();
		//transfer_timing += perf_.last();

		double dealloc_timing = Dealloc();

		perf_.store(Step(StepType::ALLOC_DEALLOC), alloc_timing + dealloc_timing /*+ transfer_timing */);

	}



	void PerformLabeling() {

		d_img_labels_.create(d_img_.size(), CV_32SC1);
		grid_size_ = dim3((d_img_.cols + BLOCK_COLS - 1) / BLOCK_COLS, (d_img_.rows + BLOCK_ROWS - 1) / BLOCK_ROWS, 1);
		block_size_ = dim3(BLOCK_COLS, BLOCK_ROWS, 1);
		
		// Phase 1
		// Etichetta i pixel localmente al blocco		
		LocalMerge << <grid_size_, block_size_ >> >(d_img_, d_img_labels_);

		// Immagine di debug della prima fase
		//cuda::GpuMat d_local_labels;
