 				Top2 left_top2;
 				left_top2.initialize();
 				for(unsigned int i = 0; i < REDUCTION_PER_THREAD; ++i){
-					left_top2.push(local_packed_cost[i]);
-				}
-				left_top2 = subgroup_merge_top2<WARP_SIZE>(left_top2);
+					left_top2.push(local_packed_cost[i], uniqueness);
+				}
+				left_top2 = subgroup_merge_top2<WARP_SIZE>(left_top2, uniqueness);
 				if(lane_id == 0){
 					left_dest[x] = compute_disparity(left_top2, uniqueness, smem_cost_sum[warp_id][smem_x]);
 				}

	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(cost1 * uniqueness >= cost0
		|| abs(disp1 - disp0) <= 1){
		int disp = disp0;
		disp <<= sgm::StereoSGM::SUBPIXEL_SHIFT;
		if (disp0 > 0 && disp0 < MAX_DISPARITY - 1) {

		const uint32_t a = __shfl_xor(x.values[0], STEP / 2, GROUP_SIZE);
		const uint32_t b = __shfl_xor(x.values[1], STEP / 2, GROUP_SIZE);
#endif
		x.push(a, uniqueness);
		x.push(b, uniqueness);
		return subgroup_merge_top2_impl<GROUP_SIZE, STEP / 2>::call(x, uniqueness);
	}
};

template <unsigned int GROUP_SIZE>
struct subgroup_merge_top2_impl<GROUP_SIZE, 1u> {
	static __device__ Top2 call(Top2 x, float){
		return x;
	}
};

template <unsigned int GROUP_SIZE>
__device__ inline Top2 subgroup_merge_top2(Top2 x, float uniqueness){
	return subgroup_merge_top2_impl<GROUP_SIZE, GROUP_SIZE>::call(x, uniqueness);
}

__device__ inline uint32_t pack_cost_index(uint32_t cost, uint32_t index){
