         int                 num_valid,          ///< [in] Number of valid elements (may be less than BLOCK_THREADS)
         ReductionOp         reduction_op)       ///< [in] Binary reduction operator
     {
-        unsigned int    warp_offset = warp_id * LOGICAL_WARP_SIZE;
-        unsigned int    warp_num_valid = (FULL_TILE && EVEN_WARP_MULTIPLE) ?
+        int         warp_offset = warp_id * LOGICAL_WARP_SIZE;
+        int         warp_num_valid = ((FULL_TILE && EVEN_WARP_MULTIPLE) || (warp_offset + LOGICAL_WARP_SIZE <= num_valid)) ?
                             LOGICAL_WARP_SIZE :
-                            (warp_offset < static_cast<unsigned int>(num_valid)) ?
-                                num_valid - warp_offset :
-                                0;
+                            num_valid - warp_offset;
 
         // Warp reduction in every warp
         T warp_aggregate = WarpReduce(temp_storage.warp_reduce[warp_id]).template Reduce<(FULL_TILE && EVEN_WARP_MULTIPLE)>(


    // Thread fields
    _TempStorage &temp_storage;
    unsigned int linear_tid;
    unsigned int warp_id;
    unsigned int lane_id;


    /// Constructor

        T                   input,          ///< [in] Calling thread's input partial reductions
        int                 num_valid)      ///< [in] Number of valid elements (may be less than BLOCK_THREADS)
    {
        cub::Sum    reduction_op;
        int         warp_offset = (warp_id * LOGICAL_WARP_SIZE);
        int         warp_num_valid = ((FULL_TILE && EVEN_WARP_MULTIPLE) || (warp_offset + LOGICAL_WARP_SIZE <= num_valid)) ?
                            LOGICAL_WARP_SIZE :
                            num_valid - warp_offset;

        // Warp reduction in every warp
        T warp_aggregate = WarpReduce(temp_storage.warp_reduce[warp_id]).template Reduce<(FULL_TILE && EVEN_WARP_MULTIPLE)>(
