device_segmented_reduce.cuh

Go to the documentation of this file.

/******************************************************************************
 * Copyright (c) 2011, Duane Merrill.  All rights reserved.
 * Copyright (c) 2011-2018, NVIDIA CORPORATION.  All rights reserved.
 *
 * Redistribution and use in source and binary forms, with or without
 * modification, are permitted provided that the following conditions are met:
 *     * Redistributions of source code must retain the above copyright
 *       notice, this list of conditions and the following disclaimer.
 *     * Redistributions in binary form must reproduce the above copyright
 *       notice, this list of conditions and the following disclaimer in the
 *       documentation and/or other materials provided with the distribution.
 *     * Neither the name of the NVIDIA CORPORATION nor the
 *       names of its contributors may be used to endorse or promote products
 *       derived from this software without specific prior written permission.
 *
 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
 * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
 * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
 * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
 * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
 * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
 * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
 * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
 * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 *
 ******************************************************************************/

#pragma once

#include <stdio.h>
#include <iterator>

#include "../iterator/arg_index_input_iterator.cuh"
#include "dispatch/dispatch_reduce.cuh"
#include "dispatch/dispatch_reduce_by_key.cuh"
#include "../util_type.cuh"
#include "../util_namespace.cuh"

CUB_NS_PREFIX

namespace cub {


struct DeviceSegmentedReduce
{
    template <
        typename            InputIteratorT,
        typename            OutputIteratorT,
        typename            OffsetIteratorT,
        typename            ReductionOp,
        typename            T>
    CUB_RUNTIME_FUNCTION
    static cudaError_t Reduce(
        void                *d_temp_storage,                    
        size_t              &temp_storage_bytes,                
        InputIteratorT      d_in,                               
        OutputIteratorT     d_out,                              
        int                 num_segments,                       
        OffsetIteratorT     d_begin_offsets,                    
        OffsetIteratorT     d_end_offsets,                      
        ReductionOp         reduction_op,                       
        T                   initial_value,                      
        cudaStream_t        stream              = 0,            
        bool                debug_synchronous   = false)        
    {
        // Signed integer type for global offsets
        typedef int OffsetT;

        return DispatchSegmentedReduce<InputIteratorT, OutputIteratorT, OffsetIteratorT, OffsetT, ReductionOp>::Dispatch(
            d_temp_storage,
            temp_storage_bytes,
            d_in,
            d_out,
            num_segments,
            d_begin_offsets,
            d_end_offsets,
            reduction_op,
            initial_value,
            stream,
            debug_synchronous);
    }


    template <
        typename            InputIteratorT,
        typename            OutputIteratorT,
        typename            OffsetIteratorT>
    CUB_RUNTIME_FUNCTION
    static cudaError_t Sum(
        void                *d_temp_storage,                    
        size_t              &temp_storage_bytes,                
        InputIteratorT      d_in,                               
        OutputIteratorT     d_out,                              
        int                 num_segments,                       
        OffsetIteratorT     d_begin_offsets,                    
        OffsetIteratorT     d_end_offsets,                      
        cudaStream_t        stream              = 0,            
        bool                debug_synchronous   = false)        
    {
        // Signed integer type for global offsets
        typedef int OffsetT;

        // The output value type
        typedef typename If<(Equals<typename std::iterator_traits<OutputIteratorT>::value_type, void>::VALUE),  // OutputT =  (if output iterator's value type is void) ?
            typename std::iterator_traits<InputIteratorT>::value_type,                                          // ... then the input iterator's value type,
            typename std::iterator_traits<OutputIteratorT>::value_type>::Type OutputT;                          // ... else the output iterator's value type

        return DispatchSegmentedReduce<InputIteratorT,  OutputIteratorT, OffsetIteratorT, OffsetT, cub::Sum>::Dispatch(
            d_temp_storage,
            temp_storage_bytes,
            d_in,
            d_out,
            num_segments,
            d_begin_offsets,
            d_end_offsets,
            cub::Sum(),
            OutputT(),            // zero-initialize
            stream,
            debug_synchronous);
    }


    template <
        typename            InputIteratorT,
        typename            OutputIteratorT,
        typename            OffsetIteratorT>
    CUB_RUNTIME_FUNCTION
    static cudaError_t Min(
        void                *d_temp_storage,                    
        size_t              &temp_storage_bytes,                
        InputIteratorT      d_in,                               
        OutputIteratorT     d_out,                              
        int                 num_segments,                       
        OffsetIteratorT     d_begin_offsets,                    
        OffsetIteratorT     d_end_offsets,                      
        cudaStream_t        stream              = 0,            
        bool                debug_synchronous   = false)        
    {
        // Signed integer type for global offsets
        typedef int OffsetT;

        // The input value type
        typedef typename std::iterator_traits<InputIteratorT>::value_type InputT;

        return DispatchSegmentedReduce<InputIteratorT,  OutputIteratorT, OffsetIteratorT, OffsetT, cub::Min>::Dispatch(
            d_temp_storage,
            temp_storage_bytes,
            d_in,
            d_out,
            num_segments,
            d_begin_offsets,
            d_end_offsets,
            cub::Min(),
            Traits<InputT>::Max(),    // replace with std::numeric_limits<T>::max() when C++11 support is more prevalent
            stream,
            debug_synchronous);
    }


    template <
        typename            InputIteratorT,
        typename            OutputIteratorT,
        typename            OffsetIteratorT>
    CUB_RUNTIME_FUNCTION
    static cudaError_t ArgMin(
        void                *d_temp_storage,                    
        size_t              &temp_storage_bytes,                
        InputIteratorT      d_in,                               
        OutputIteratorT     d_out,                              
        int                 num_segments,                       
        OffsetIteratorT     d_begin_offsets,                    
        OffsetIteratorT     d_end_offsets,                      
        cudaStream_t        stream              = 0,            
        bool                debug_synchronous   = false)        
    {
        // Signed integer type for global offsets
        typedef int OffsetT;

        // The input type
        typedef typename std::iterator_traits<InputIteratorT>::value_type InputValueT;

        // The output tuple type
        typedef typename If<(Equals<typename std::iterator_traits<OutputIteratorT>::value_type, void>::VALUE),  // OutputT =  (if output iterator's value type is void) ?
            KeyValuePair<OffsetT, InputValueT>,                                                                 // ... then the key value pair OffsetT + InputValueT
            typename std::iterator_traits<OutputIteratorT>::value_type>::Type OutputTupleT;                     // ... else the output iterator's value type

        // The output value type
        typedef typename OutputTupleT::Value OutputValueT;

        // Wrapped input iterator to produce index-value <OffsetT, InputT> tuples
        typedef ArgIndexInputIterator<InputIteratorT, OffsetT, OutputValueT> ArgIndexInputIteratorT;
        ArgIndexInputIteratorT d_indexed_in(d_in);

        // Initial value
        OutputTupleT initial_value(1, Traits<InputValueT>::Max());   // replace with std::numeric_limits<T>::max() when C++11 support is more prevalent

        return DispatchSegmentedReduce<ArgIndexInputIteratorT,  OutputIteratorT, OffsetIteratorT, OffsetT, cub::ArgMin>::Dispatch(
            d_temp_storage,
            temp_storage_bytes,
            d_indexed_in,
            d_out,
            num_segments,
            d_begin_offsets,
            d_end_offsets,
            cub::ArgMin(),
            initial_value,
            stream,
            debug_synchronous);
    }


    template <
        typename            InputIteratorT,
        typename            OutputIteratorT,
        typename            OffsetIteratorT>
    CUB_RUNTIME_FUNCTION
    static cudaError_t Max(
        void                *d_temp_storage,                    
        size_t              &temp_storage_bytes,                
        InputIteratorT      d_in,                               
        OutputIteratorT     d_out,                              
        int                 num_segments,                       
        OffsetIteratorT     d_begin_offsets,                    
        OffsetIteratorT     d_end_offsets,                      
        cudaStream_t        stream              = 0,            
        bool                debug_synchronous   = false)        
    {
        // Signed integer type for global offsets
        typedef int OffsetT;

        // The input value type
        typedef typename std::iterator_traits<InputIteratorT>::value_type InputT;

        return DispatchSegmentedReduce<InputIteratorT,  OutputIteratorT, OffsetIteratorT, OffsetT, cub::Max>::Dispatch(
            d_temp_storage,
            temp_storage_bytes,
            d_in,
            d_out,
            num_segments,
            d_begin_offsets,
            d_end_offsets,
            cub::Max(),
            Traits<InputT>::Lowest(),    // replace with std::numeric_limits<T>::lowest() when C++11 support is more prevalent
            stream,
            debug_synchronous);
    }


    template <
        typename            InputIteratorT,
        typename            OutputIteratorT,
        typename            OffsetIteratorT>
    CUB_RUNTIME_FUNCTION
    static cudaError_t ArgMax(
        void                *d_temp_storage,                    
        size_t              &temp_storage_bytes,                
        InputIteratorT      d_in,                               
        OutputIteratorT     d_out,                              
        int                 num_segments,                       
        OffsetIteratorT     d_begin_offsets,                    
        OffsetIteratorT     d_end_offsets,                      
        cudaStream_t        stream              = 0,            
        bool                debug_synchronous   = false)        
    {
        // Signed integer type for global offsets
        typedef int OffsetT;

        // The input type
        typedef typename std::iterator_traits<InputIteratorT>::value_type InputValueT;

        // The output tuple type
        typedef typename If<(Equals<typename std::iterator_traits<OutputIteratorT>::value_type, void>::VALUE),  // OutputT =  (if output iterator's value type is void) ?
            KeyValuePair<OffsetT, InputValueT>,                                                                 // ... then the key value pair OffsetT + InputValueT
            typename std::iterator_traits<OutputIteratorT>::value_type>::Type OutputTupleT;                     // ... else the output iterator's value type

        // The output value type
        typedef typename OutputTupleT::Value OutputValueT;

        // Wrapped input iterator to produce index-value <OffsetT, InputT> tuples
        typedef ArgIndexInputIterator<InputIteratorT, OffsetT, OutputValueT> ArgIndexInputIteratorT;
        ArgIndexInputIteratorT d_indexed_in(d_in);

        // Initial value
        OutputTupleT initial_value(1, Traits<InputValueT>::Lowest());     // replace with std::numeric_limits<T>::lowest() when C++11 support is more prevalent

        return DispatchSegmentedReduce<ArgIndexInputIteratorT, OutputIteratorT, OffsetIteratorT, OffsetT, cub::ArgMax>::Dispatch(
            d_temp_storage,
            temp_storage_bytes,
            d_indexed_in,
            d_out,
            num_segments,
            d_begin_offsets,
            d_end_offsets,
            cub::ArgMax(),
            initial_value,
            stream,
            debug_synchronous);
    }

};

}               // CUB namespace
CUB_NS_POSTFIX  // Optional outer namespace(s)