 #include "../../util_type.cuh"
 
 #include <stdint.h>
+#include <type_traits>
+
+#include <nv/target>
 
 CUB_NAMESPACE_BEGIN
 

        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;
    }



    //---------------------------------------------------------------------
    // 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 <
