 thrust::host_vector<sgm::output_type> winner_takes_all_left(
 	const thrust::host_vector<sgm::cost_type>& src,
 	size_t width, size_t height, size_t disparity,
-	float uniqueness)
+	float uniqueness, bool subpixel)
 {
 	thrust::host_vector<sgm::output_type> result(width * height);
 	for(size_t i = 0; i < height; ++i){


}


TEST(WinnerTakesAllTest, RandomLeft){
	static constexpr size_t width = 313, height = 237, disparity = 128;
	static constexpr float uniqueness = 0.95f;
	const auto input = generate_random_sequence<sgm::cost_type>(
		width * height * disparity * NUM_PATHS);
	const auto expect = winner_takes_all_left(
		input, width, height, disparity, uniqueness);

	sgm::WinnerTakesAll<disparity> wta;
	const auto d_input = to_device_vector(input);
	wta.enqueue(d_input.data().get(), width, height, uniqueness, 0);
	cudaStreamSynchronize(0);

	const thrust::device_vector<sgm::output_type> d_actual(
		wta.get_left_output(), wta.get_left_output() + (width * height));
	const auto actual = to_host_vector(d_actual);
	EXPECT_EQ(actual, expect);
	debug_compare(actual.data(), expect.data(), width, height, 1);
}

TEST(WinnerTakesAllTest, RandomRight){

				const int cost1 = v[1].first;
				const int disp0 = v[0].second;
				const int disp1 = v[1].second;
				sgm::output_type dst;
				if (cost1 * uniqueness < cost0 && abs(disp0 - disp1) > 1) {
					dst = 0;
				} else {
					dst = disp0;
					if (subpixel) {
						dst <<= sgm::StereoSGM::SUBPIXEL_SHIFT;
						if (0 < disp0 && disp0 < disparity - 1) {
							const int left = w[disp0 - 1].first;
							const int right = w[disp0 + 1].first;
							const int numer = left - right;
							const int denom = left - 2 * cost0 + right;
							dst += ((numer << sgm::StereoSGM::SUBPIXEL_SHIFT) + denom) / (2 * denom);
						}
					}
				}
				result[i * width + j] = dst;
			}
		}
	}
