 		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);
-		x.push(b);
-		return subgroup_merge_top2_impl<GROUP_SIZE, STEP / 2>::call(x);
+		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){
+	static __device__ Top2 call(Top2 x, float){
 		return x;
 	}
 };
 
 template <unsigned int GROUP_SIZE>
-__device__ inline Top2 subgroup_merge_top2(Top2 x){
-	return subgroup_merge_top2_impl<GROUP_SIZE, GROUP_SIZE>::call(x);
+__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){

static constexpr unsigned int BLOCK_SIZE = WARPS_PER_BLOCK * WARP_SIZE;


__device__ inline void update_top2(uint32_t& v0, uint32_t& v1, uint32_t x){
	const uint32_t y = max(x, v0);
	v0 = min(x, v0);
	v1 = min(y, v1);
}

struct Top2 {

				Top2 left_top2;
				left_top2.initialize();
				for(unsigned int i = 0; i < REDUCTION_PER_THREAD; ++i){
					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]);
				}
