 					right_top2[i].push(recv);
 					if(d == MAX_DISPARITY - 1){
 						if(0 <= p){
-							right_dest[p] = compute_disparity(right_top2[i], uniqueness);
+							right_dest[p] = compute_disparity_normal<MAX_DISPARITY>(right_top2[i], uniqueness);
 						}
 						right_top2[i].initialize();
 					}

	return packed & 0xffffu;
}

__device__ inline uint32_t compute_disparity(Top2 t2, float uniqueness){
	const float cost0 = static_cast<float>(unpack_cost(t2.values[0]));
	const float cost1 = static_cast<float>(unpack_cost(t2.values[1]));
	const int disp0 = static_cast<int>(unpack_index(t2.values[0]));
	const int disp1 = static_cast<int>(unpack_index(t2.values[1]));
	if(static_cast<float>(cost1) * uniqueness >= static_cast<float>(cost0)){
		return static_cast<uint32_t>(disp0);
	}else if(abs(disp1 - disp0) <= 1){
		return static_cast<uint32_t>(disp0);
	}else{
		return 0;
	}
}

template <unsigned int MAX_DISPARITY>
__global__ void winner_takes_all_kernel(
	output_type *left_dest,
	output_type *right_dest,

		const unsigned int k = lane_id * REDUCTION_PER_THREAD + i;
		const int p = static_cast<int>(((width - k) & ~(MAX_DISPARITY - 1)) + k);
		if(p < width){
			right_dest[p] = compute_disparity_normal<MAX_DISPARITY>(right_top2[i], uniqueness);
		}
	}
}
