/// Computes a threadblock-wide reduction using addition (+) as the reduction operator. The first num_valid threads each contribute one reduction partial. The return value is only valid for thread<sub>0</sub>.
template <bool FULL_TILE>
__device__ __forceinline__ T Sum(
T partial, ///< [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;
if (WARP_SYNCHRONOUS)
{
// Short-circuit directly to warp synchronous reduction (unguarded if active threads is a power-of-two)
/// Computes a threadblock-wide reduction using the specified reduction operator. The first num_valid threads each contribute one reduction partial. The return value is only valid for thread<sub>0</sub>.
template <
bool FULL_TILE,
typename ReductionOp>
__device__ __forceinline__ T Reduce(
T partial, ///< [in] Calling thread's input partial reductions
int num_valid, ///< [in] Number of valid elements (may be less than BLOCK_THREADS)