device_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 <limits>
 
#include "../iterator/arg_index_input_iterator.cuh"
#include "dispatch/dispatch_reduce.cuh"
#include "dispatch/dispatch_reduce_by_key.cuh"
#include "../util_namespace.cuh"
 
CUB_NS_PREFIX
 
namespace cub {
 
 
struct DeviceReduce
{
    template <
        typename                    InputIteratorT,
        typename                    OutputIteratorT,
        typename                    ReductionOpT,
        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_items,                          
        ReductionOpT                reduction_op,                       
        T                           init,                               
        cudaStream_t                stream              = 0,            
        bool                        debug_synchronous   = false)        
    {
        // Signed integer type for global offsets
        typedef int OffsetT;
 
        return DispatchReduce<InputIteratorT, OutputIteratorT, OffsetT, ReductionOpT>::Dispatch(
            d_temp_storage,
            temp_storage_bytes,
            d_in,
            d_out,
            num_items,
            reduction_op,
            init,
            stream,
            debug_synchronous);
    }
 
 
    template <
        typename                    InputIteratorT,
        typename                    OutputIteratorT>
    CUB_RUNTIME_FUNCTION
    static cudaError_t Sum(
        void                        *d_temp_storage,                    
        size_t                      &temp_storage_bytes,                
        InputIteratorT              d_in,                               
        OutputIteratorT             d_out,                              
        int                         num_items,                          
        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 DispatchReduce<InputIteratorT, OutputIteratorT, OffsetT, cub::Sum>::Dispatch(
            d_temp_storage,
            temp_storage_bytes,
            d_in,
            d_out,
            num_items,
            cub::Sum(),
            OutputT(),            // zero-initialize
            stream,
            debug_synchronous);
    }
 
 
    template <
        typename                    InputIteratorT,
        typename                    OutputIteratorT>
    CUB_RUNTIME_FUNCTION
    static cudaError_t Min(
        void                        *d_temp_storage,                    
        size_t                      &temp_storage_bytes,                
        InputIteratorT              d_in,                               
        OutputIteratorT             d_out,                              
        int                         num_items,                          
        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 DispatchReduce<InputIteratorT, OutputIteratorT, OffsetT, cub::Min>::Dispatch(
            d_temp_storage,
            temp_storage_bytes,
            d_in,
            d_out,
            num_items,
            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>
    CUB_RUNTIME_FUNCTION
    static cudaError_t ArgMin(
        void                        *d_temp_storage,                    
        size_t                      &temp_storage_bytes,                
        InputIteratorT              d_in,                               
        OutputIteratorT             d_out,                              
        int                         num_items,                          
        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 DispatchReduce<ArgIndexInputIteratorT, OutputIteratorT, OffsetT, cub::ArgMin>::Dispatch(
            d_temp_storage,
            temp_storage_bytes,
            d_indexed_in,
            d_out,
            num_items,
            cub::ArgMin(),
            initial_value,
            stream,
            debug_synchronous);
    }
 
 
    template <
        typename                    InputIteratorT,
        typename                    OutputIteratorT>
    CUB_RUNTIME_FUNCTION
    static cudaError_t Max(
        void                        *d_temp_storage,                    
        size_t                      &temp_storage_bytes,                
        InputIteratorT              d_in,                               
        OutputIteratorT             d_out,                              
        int                         num_items,                          
        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 DispatchReduce<InputIteratorT, OutputIteratorT, OffsetT, cub::Max>::Dispatch(
            d_temp_storage,
            temp_storage_bytes,
            d_in,
            d_out,
            num_items,
            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>
    CUB_RUNTIME_FUNCTION
    static cudaError_t ArgMax(
        void                        *d_temp_storage,                    
        size_t                      &temp_storage_bytes,                
        InputIteratorT              d_in,                               
        OutputIteratorT             d_out,                              
        int                         num_items,                          
        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 DispatchReduce<ArgIndexInputIteratorT, OutputIteratorT, OffsetT, cub::ArgMax>::Dispatch(
            d_temp_storage,
            temp_storage_bytes,
            d_indexed_in,
            d_out,
            num_items,
            cub::ArgMax(),
            initial_value,
            stream,
            debug_synchronous);
    }
 
 
    template <
        typename                    KeysInputIteratorT,
        typename                    UniqueOutputIteratorT,
        typename                    ValuesInputIteratorT,
        typename                    AggregatesOutputIteratorT,
        typename                    NumRunsOutputIteratorT,
        typename                    ReductionOpT>
    CUB_RUNTIME_FUNCTION __forceinline__
    static cudaError_t ReduceByKey(
        void                        *d_temp_storage,                
        size_t                      &temp_storage_bytes,            
        KeysInputIteratorT          d_keys_in,                      
        UniqueOutputIteratorT       d_unique_out,                   
        ValuesInputIteratorT        d_values_in,                    
        AggregatesOutputIteratorT   d_aggregates_out,               
        NumRunsOutputIteratorT      d_num_runs_out,                 
        ReductionOpT                reduction_op,                   
        int                         num_items,                      
        cudaStream_t                stream             = 0,         
        bool                        debug_synchronous  = false)     
    {
        // Signed integer type for global offsets
        typedef int OffsetT;
 
        // FlagT iterator type (not used)
 
        // Selection op (not used)
 
        // Default == operator
        typedef Equality EqualityOp;
 
        return DispatchReduceByKey<KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOp, ReductionOpT, OffsetT>::Dispatch(
            d_temp_storage,
            temp_storage_bytes,
            d_keys_in,
            d_unique_out,
            d_values_in,
            d_aggregates_out,
            d_num_runs_out,
            EqualityOp(),
            reduction_op,
            num_items,
            stream,
            debug_synchronous);
    }
 
};
 
}               // CUB namespace
CUB_NS_POSTFIX  // Optional outer namespace(s)