     //---------------------------------------------------------------------
     // Reduction operations
     //---------------------------------------------------------------------
+    template <typename ReductionOp>
+    __device__ __forceinline__ T ReduceImpl(
+        Int2Type<0>     /* all_lanes_valid */, 
+        T               input,                  ///< [in] Calling thread's input
+        int             valid_items,            ///< [in] Total number of valid items across the logical warp
+        ReductionOp     reduction_op)           ///< [in] Binary reduction operator
+    {
+        int last_lane = valid_items - 1;
+
+        T output = input;
+
+        // Template-iterate reduction steps
+        ReduceStep(output, reduction_op, last_lane, Int2Type<0>());
+
+        return output;
+    }
+
+    template <typename ReductionOp>
+    __device__ __forceinline__ T ReduceImpl(
+        Int2Type<1>     /* all_lanes_valid */, 
+        T               input,                  ///< [in] Calling thread's input
+        int             /* valid_items */,      ///< [in] Total number of valid items across the logical warp
+        ReductionOp     reduction_op)           ///< [in] Binary reduction operator
+    {
+        int last_lane = LOGICAL_WARP_THREADS - 1;
+
+        T output = input;
+
+        // Template-iterate reduction steps
+        ReduceStep(output, reduction_op, last_lane, Int2Type<0>());
+
+        return output;
+    }
+
+    template <class U = T>
+    __device__ __forceinline__ 
+    typename std::enable_if<
+               std::is_same<int, U>::value 
+            || std::is_same<unsigned int, U>::value, T>::type
+    ReduceImpl(Int2Type<1> /* all_lanes_valid */,
+               T input,
+               int /* valid_items */,
+               cub::Sum /* reduction_op */)
+    {
+      T output = input;
+
+      NV_IF_TARGET(NV_PROVIDES_SM_80,
+                   (output = __reduce_add_sync(member_mask, input);),
+                   (output = ReduceImpl<cub::Sum>(Int2Type<1>{},
+                                                  input,
+                                                  LOGICAL_WARP_THREADS,
+                                                  cub::Sum{});));
+
+      return output;
+    }
+
+    template <class U = T>
+    __device__ __forceinline__ 
+    typename std::enable_if<
+               std::is_same<int, U>::value 
+            || std::is_same<unsigned int, U>::value, T>::type
+    ReduceImpl(Int2Type<1> /* all_lanes_valid */,
+               T input,
+               int /* valid_items */,
+               cub::Min /* reduction_op */)
+    {
+      T output = input;
+
+      NV_IF_TARGET(NV_PROVIDES_SM_80,
+                   (output = __reduce_min_sync(member_mask, input);),
+                   (output = ReduceImpl<cub::Min>(Int2Type<1>{},
+                                                  input,
+                                                  LOGICAL_WARP_THREADS,
+                                                  cub::Min{});));
+
+      return output;
+    }
+
+    template <class U = T>
+    __device__ __forceinline__ 
+    typename std::enable_if<
+               std::is_same<int, U>::value 
+            || std::is_same<unsigned int, U>::value, T>::type
+    ReduceImpl(Int2Type<1> /* all_lanes_valid */,
+               T input,
+               int /* valid_items */,
+               cub::Max /* reduction_op */)
+    {
+      T output = input;
+
+      NV_IF_TARGET(NV_PROVIDES_SM_80,
+                   (output = __reduce_max_sync(member_mask, input);),
+                   (output = ReduceImpl<cub::Max>(Int2Type<1>{},
+                                                  input,
+                                                  LOGICAL_WARP_THREADS,
+                                                  cub::Max{});));
+
+      return output;
+    }
 
     /// Reduction
     template <

        int             valid_items,            ///< [in] Total number of valid items across the logical warp
        ReductionOp     reduction_op)           ///< [in] Binary reduction operator
    {
        int last_lane = (ALL_LANES_VALID) ?
                            LOGICAL_WARP_THREADS - 1 :
                            valid_items - 1;

        T output = input;

//        // Iterate reduction steps
//        #pragma unroll
//        for (int STEP = 0; STEP < STEPS; STEP++)
//        {
//            output = ReduceStep(output, reduction_op, last_lane, 1 << STEP, Int2Type<IsInteger<T>::IS_SMALL_UNSIGNED>());
//        }

        // Template-iterate reduction steps
        ReduceStep(output, reduction_op, last_lane, Int2Type<0>());

        return output;
    }



#include "../../util_type.cuh"

#include <stdint.h>
#include <type_traits>

#include <nv/target>

CUB_NAMESPACE_BEGIN

