 		values[1] = 0xffffffffu;
 	}
 
-	__device__ void push(uint32_t x){
-		update_top2(values[0], values[1], x);
+	__device__ void push(uint32_t x, float uniqueness){
+		update_top2(values[0], values[1], x, uniqueness);
 	}
 };
 
 template <unsigned int GROUP_SIZE, unsigned int STEP>
 struct subgroup_merge_top2_impl {
-	static __device__ Top2 call(Top2 x){
+	static __device__ Top2 call(Top2 x, float uniqueness){
 #if CUDA_VERSION >= 9000
 		const uint32_t a = __shfl_xor_sync(0xffffffffu, x.values[0], STEP / 2, GROUP_SIZE);
 		const uint32_t b = __shfl_xor_sync(0xffffffffu, x.values[1], STEP / 2, GROUP_SIZE);

		uint32_t uint32;
		ushort2 uint16x2;
	} u;
	u.uint16x2.x = static_cast<uint16_t>(index);
	u.uint16x2.y = static_cast<uint16_t>(cost);
	return u.uint32;
}

__device__ uint32_t unpack_cost(uint32_t packed){
	return packed >> 16;
}

__device__ uint32_t unpack_index(uint32_t packed){
	return packed & 0xffffu;
}

using ComputeDisparity = uint32_t(*)(Top2, float, uint16_t*);

	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(unpack_uniq(t2.values[0])){
		return disp0;
	}else{
		return 0;
