             static_assert(LOGICAL_WARP_THREADS && !(LOGICAL_WARP_THREADS & (LOGICAL_WARP_THREADS-1)), "Logical warp size must be a power of 2");
             }
 
+        //! Sum reduction.
+        /*!
+         * \param input Thread data to sum.
+         * \returns output The result of the sum reduction in thread 0 of the (sub-)warp.
+         *
+         * The sum reduction for a 4-thread sub-warp with \a input <tt>{1,2,3,4}</tt> gives \a output <tt>10</tt> in thread 0.
+         * The result in the other threads is undefined.
+         */
         DEVICE T Sum(T input)
             {
             return Reduce(input, cub::Sum());
             }
 
+        //! Sum reduction over valid items.
+        /*!
+         * \param input Thread data to sum.
+         * \param valid_items Total number of valid items in the (sub)-warp.
+         * \returns output The result of the sum reduction in thread 0 of the (sub-)warp.
+         *
+         * The number of valid items may be smaller than the (sub-)warp. For example, if \a valid items is 3, then
+         * the sum reduction for a 4-thread sub-warp with \a input <tt>{1,2,3,4}</tt> gives \a output <tt>6</tt> in thread 0.
+         * The result in the other threads is undefined.
+         */
         DEVICE T Sum(T input, int valid_items)
             {
             return Reduce(input, cub::Sum(), valid_items);
             }
 
+        //! Custom reduction.
+        /*!
+         * \param input Thread data to sum.
+         * \param reduce_op Custom reduction operation.
+         * \returns output The result of the reduction in thread 0 of the (sub-)warp.
+         *
+         * \tparam ReduceOpT The type of the reduction operation.
+         *
+         * This is a generalization of Sum() to custom operators.
+         */
         template<typename ReduceOpT>
         DEVICE T Reduce(T input, ReduceOpT reduce_op)
             {

namespace detail
{

template<typename T, int LOGICAL_WARP_THREADS = CUB_PTX_WARP_THREADS, int PTX_ARCH = CUB_PTX_ARCH>
class WarpReduce
    {

            return WarpReduceShfl(tmp).template Reduce<true>(input, LOGICAL_WARP_THREADS, reduce_op);
            }

        //! Custom reduction over valid items.
        /*!
         * \param input Thread data to sum.
         * \param reduce_op Custom reduction operation.
         * \param valid_items Total number of valid items in the (sub)-warp.
         * \returns output The result of the reduction in thread 0 of the (sub-)warp.
         *
         * \tparam ReduceOpT The type of the reduction operation.
         *
         * This is a generalization of Sum() over valid items to custom operators.
         */
        template<typename ReduceOpT>
        DEVICE T Reduce(T input, ReduceOpT reduce_op, int valid_items)
            {
