 				const int cost1 = v[1].first;
 				const int disp0 = v[0].second;
 				const int disp1 = v[1].second;
-				result[i * width + j] = static_cast<sgm::output_type>(
-					(cost1 * uniqueness < cost0 && abs(disp0 - disp1) > 1)
-						? 0
-						: disp0);
+				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;
 			}
 		}
 	}


	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(

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, bool subpixel)
{
	thrust::host_vector<sgm::output_type> result(width * height);
	for(size_t i = 0; i < height; ++i){
