 private:
 	void Alloc() {
 		d_img_labels_.create(d_img_.size(), CV_32SC1);
+		if ((d_img_.rows == 1 || d_img_.cols == 1) && !((d_img_.rows + d_img_.cols) % 2)) {
+			cudaMalloc(&last_pixel_, sizeof(unsigned char));
+			last_pixel_allocated_ = true;
+		}
+		else {
+			// last_pixel_ = d_img_labels_.data + d_img_labels_.step + sizeof(unsigned int);
+			last_pixel_ = d_img_labels_.data + ((d_img_labels_.rows - 2) * d_img_labels_.step) + (d_img_labels_.cols - 2) * d_img_labels_.elemSize();
+		}
 	}
 
 	void Dealloc() {
+		if (last_pixel_allocated_) {
+			cudaFree(last_pixel_);
+		}
 	}
 
 	double MemoryTransferHostToDevice() {


		d_img_labels_.create(d_img_.size(), CV_32SC1);

		grid_size_ = dim3((((d_img_.cols + 1) / 2) + BLOCK_COLS - 1) / BLOCK_COLS, (((d_img_.rows + 1) / 2) + BLOCK_ROWS - 1) / BLOCK_ROWS, 1);
		block_size_ = dim3(BLOCK_COLS, BLOCK_ROWS, 1);

		InitLabeling << <grid_size_, block_size_ >> > (d_img_, d_img_labels_);

		Compression << <grid_size_, block_size_ >> > (d_img_labels_);

		//Mat1i init_labels;
		//d_img_labels_.download(init_labels);

		//cuda::GpuMat d_expanded_connections;
		//d_expanded_connections.create(d_connections_.rows * 3, d_connections_.cols * 3, CV_8UC1);

private:
	dim3 grid_size_;
	dim3 block_size_;
	unsigned char *last_pixel_;
	bool last_pixel_allocated_;

public:
	CUDA_BKE() {}
