Spaces:
Runtime error
Runtime error
| /****************************************************************************** | |
| * 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. | |
| * | |
| ******************************************************************************/ | |
| /** | |
| * \file | |
| * cub::DeviceReduce provides device-wide, parallel operations for computing a reduction across a sequence of data items residing within device-accessible memory. | |
| */ | |
| #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 "../config.cuh" | |
| /// Optional outer namespace(s) | |
| CUB_NS_PREFIX | |
| /// CUB namespace | |
| namespace cub { | |
| /** | |
| * \brief DeviceReduce provides device-wide, parallel operations for computing a reduction across a sequence of data items residing within device-accessible memory.  | |
| * \ingroup SingleModule | |
| * | |
| * \par Overview | |
| * A <a href="http://en.wikipedia.org/wiki/Reduce_(higher-order_function)"><em>reduction</em></a> (or <em>fold</em>) | |
| * uses a binary combining operator to compute a single aggregate from a sequence of input elements. | |
| * | |
| * \par Usage Considerations | |
| * \cdp_class{DeviceReduce} | |
| * | |
| * \par Performance | |
| * \linear_performance{reduction, reduce-by-key, and run-length encode} | |
| * | |
| * \par | |
| * The following chart illustrates DeviceReduce::Sum | |
| * performance across different CUDA architectures for \p int32 keys. | |
| * | |
| * \image html reduce_int32.png | |
| * | |
| * \par | |
| * The following chart illustrates DeviceReduce::ReduceByKey (summation) | |
| * performance across different CUDA architectures for \p fp32 | |
| * values. Segments are identified by \p int32 keys, and have lengths uniformly sampled from [1,1000]. | |
| * | |
| * \image html reduce_by_key_fp32_len_500.png | |
| * | |
| * \par | |
| * \plots_below | |
| * | |
| */ | |
| struct DeviceReduce | |
| { | |
| /** | |
| * \brief Computes a device-wide reduction using the specified binary \p reduction_op functor and initial value \p init. | |
| * | |
| * \par | |
| * - Does not support binary reduction operators that are non-commutative. | |
| * - Provides "run-to-run" determinism for pseudo-associative reduction | |
| * (e.g., addition of floating point types) on the same GPU device. | |
| * However, results for pseudo-associative reduction may be inconsistent | |
| * from one device to a another device of a different compute-capability | |
| * because CUB can employ different tile-sizing for different architectures. | |
| * - \devicestorage | |
| * | |
| * \par Snippet | |
| * The code snippet below illustrates a user-defined min-reduction of a device vector of \p int data elements. | |
| * \par | |
| * \code | |
| * #include <cub/cub.cuh> // or equivalently <cub/device/device_radix_sort.cuh> | |
| * | |
| * // CustomMin functor | |
| * struct CustomMin | |
| * { | |
| * template <typename T> | |
| * __device__ __forceinline__ | |
| * T operator()(const T &a, const T &b) const { | |
| * return (b < a) ? b : a; | |
| * } | |
| * }; | |
| * | |
| * // Declare, allocate, and initialize device-accessible pointers for input and output | |
| * int num_items; // e.g., 7 | |
| * int *d_in; // e.g., [8, 6, 7, 5, 3, 0, 9] | |
| * int *d_out; // e.g., [-] | |
| * CustomMin min_op; | |
| * int init; // e.g., INT_MAX | |
| * ... | |
| * | |
| * // Determine temporary device storage requirements | |
| * void *d_temp_storage = NULL; | |
| * size_t temp_storage_bytes = 0; | |
| * cub::DeviceReduce::Reduce(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items, min_op, init); | |
| * | |
| * // Allocate temporary storage | |
| * cudaMalloc(&d_temp_storage, temp_storage_bytes); | |
| * | |
| * // Run reduction | |
| * cub::DeviceReduce::Reduce(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items, min_op, init); | |
| * | |
| * // d_out <-- [0] | |
| * | |
| * \endcode | |
| * | |
| * \tparam InputIteratorT <b>[inferred]</b> Random-access input iterator type for reading input items \iterator | |
| * \tparam OutputIteratorT <b>[inferred]</b> Output iterator type for recording the reduced aggregate \iterator | |
| * \tparam ReductionOpT <b>[inferred]</b> Binary reduction functor type having member <tt>T operator()(const T &a, const T &b)</tt> | |
| * \tparam T <b>[inferred]</b> Data element type that is convertible to the \p value type of \p InputIteratorT | |
| */ | |
| template < | |
| typename InputIteratorT, | |
| typename OutputIteratorT, | |
| typename ReductionOpT, | |
| typename T> | |
| CUB_RUNTIME_FUNCTION | |
| static cudaError_t Reduce( | |
| void *d_temp_storage, ///< [in] %Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done. | |
| size_t &temp_storage_bytes, ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation | |
| InputIteratorT d_in, ///< [in] Pointer to the input sequence of data items | |
| OutputIteratorT d_out, ///< [out] Pointer to the output aggregate | |
| int num_items, ///< [in] Total number of input items (i.e., length of \p d_in) | |
| ReductionOpT reduction_op, ///< [in] Binary reduction functor | |
| T init, ///< [in] Initial value of the reduction | |
| cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>. | |
| bool debug_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Also causes launch configurations to be printed to the console. Default is \p 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); | |
| } | |
| /** | |
| * \brief Computes a device-wide sum using the addition (\p +) operator. | |
| * | |
| * \par | |
| * - Uses \p 0 as the initial value of the reduction. | |
| * - Does not support \p + operators that are non-commutative.. | |
| * - Provides "run-to-run" determinism for pseudo-associative reduction | |
| * (e.g., addition of floating point types) on the same GPU device. | |
| * However, results for pseudo-associative reduction may be inconsistent | |
| * from one device to a another device of a different compute-capability | |
| * because CUB can employ different tile-sizing for different architectures. | |
| * - \devicestorage | |
| * | |
| * \par Performance | |
| * The following charts illustrate saturated sum-reduction performance across different | |
| * CUDA architectures for \p int32 and \p int64 items, respectively. | |
| * | |
| * \image html reduce_int32.png | |
| * \image html reduce_int64.png | |
| * | |
| * \par Snippet | |
| * The code snippet below illustrates the sum-reduction of a device vector of \p int data elements. | |
| * \par | |
| * \code | |
| * #include <cub/cub.cuh> // or equivalently <cub/device/device_radix_sort.cuh> | |
| * | |
| * // Declare, allocate, and initialize device-accessible pointers for input and output | |
| * int num_items; // e.g., 7 | |
| * int *d_in; // e.g., [8, 6, 7, 5, 3, 0, 9] | |
| * int *d_out; // e.g., [-] | |
| * ... | |
| * | |
| * // Determine temporary device storage requirements | |
| * void *d_temp_storage = NULL; | |
| * size_t temp_storage_bytes = 0; | |
| * cub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items); | |
| * | |
| * // Allocate temporary storage | |
| * cudaMalloc(&d_temp_storage, temp_storage_bytes); | |
| * | |
| * // Run sum-reduction | |
| * cub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items); | |
| * | |
| * // d_out <-- [38] | |
| * | |
| * \endcode | |
| * | |
| * \tparam InputIteratorT <b>[inferred]</b> Random-access input iterator type for reading input items \iterator | |
| * \tparam OutputIteratorT <b>[inferred]</b> Output iterator type for recording the reduced aggregate \iterator | |
| */ | |
| template < | |
| typename InputIteratorT, | |
| typename OutputIteratorT> | |
| CUB_RUNTIME_FUNCTION | |
| static cudaError_t Sum( | |
| void *d_temp_storage, ///< [in] %Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done. | |
| size_t &temp_storage_bytes, ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation | |
| InputIteratorT d_in, ///< [in] Pointer to the input sequence of data items | |
| OutputIteratorT d_out, ///< [out] Pointer to the output aggregate | |
| int num_items, ///< [in] Total number of input items (i.e., length of \p d_in) | |
| cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>. | |
| bool debug_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Also causes launch configurations to be printed to the console. Default is \p 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); | |
| } | |
| /** | |
| * \brief Computes a device-wide minimum using the less-than ('<') operator. | |
| * | |
| * \par | |
| * - Uses <tt>std::numeric_limits<T>::max()</tt> as the initial value of the reduction. | |
| * - Does not support \p < operators that are non-commutative. | |
| * - Provides "run-to-run" determinism for pseudo-associative reduction | |
| * (e.g., addition of floating point types) on the same GPU device. | |
| * However, results for pseudo-associative reduction may be inconsistent | |
| * from one device to a another device of a different compute-capability | |
| * because CUB can employ different tile-sizing for different architectures. | |
| * - \devicestorage | |
| * | |
| * \par Snippet | |
| * The code snippet below illustrates the min-reduction of a device vector of \p int data elements. | |
| * \par | |
| * \code | |
| * #include <cub/cub.cuh> // or equivalently <cub/device/device_radix_sort.cuh> | |
| * | |
| * // Declare, allocate, and initialize device-accessible pointers for input and output | |
| * int num_items; // e.g., 7 | |
| * int *d_in; // e.g., [8, 6, 7, 5, 3, 0, 9] | |
| * int *d_out; // e.g., [-] | |
| * ... | |
| * | |
| * // Determine temporary device storage requirements | |
| * void *d_temp_storage = NULL; | |
| * size_t temp_storage_bytes = 0; | |
| * cub::DeviceReduce::Min(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items); | |
| * | |
| * // Allocate temporary storage | |
| * cudaMalloc(&d_temp_storage, temp_storage_bytes); | |
| * | |
| * // Run min-reduction | |
| * cub::DeviceReduce::Min(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items); | |
| * | |
| * // d_out <-- [0] | |
| * | |
| * \endcode | |
| * | |
| * \tparam InputIteratorT <b>[inferred]</b> Random-access input iterator type for reading input items \iterator | |
| * \tparam OutputIteratorT <b>[inferred]</b> Output iterator type for recording the reduced aggregate \iterator | |
| */ | |
| template < | |
| typename InputIteratorT, | |
| typename OutputIteratorT> | |
| CUB_RUNTIME_FUNCTION | |
| static cudaError_t Min( | |
| void *d_temp_storage, ///< [in] %Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done. | |
| size_t &temp_storage_bytes, ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation | |
| InputIteratorT d_in, ///< [in] Pointer to the input sequence of data items | |
| OutputIteratorT d_out, ///< [out] Pointer to the output aggregate | |
| int num_items, ///< [in] Total number of input items (i.e., length of \p d_in) | |
| cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>. | |
| bool debug_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Also causes launch configurations to be printed to the console. Default is \p 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); | |
| } | |
| /** | |
| * \brief Finds the first device-wide minimum using the less-than ('<') operator, also returning the index of that item. | |
| * | |
| * \par | |
| * - The output value type of \p d_out is cub::KeyValuePair <tt><int, T></tt> (assuming the value type of \p d_in is \p T) | |
| * - The minimum is written to <tt>d_out.value</tt> and its offset in the input array is written to <tt>d_out.key</tt>. | |
| * - The <tt>{1, std::numeric_limits<T>::max()}</tt> tuple is produced for zero-length inputs | |
| * - Does not support \p < operators that are non-commutative. | |
| * - Provides "run-to-run" determinism for pseudo-associative reduction | |
| * (e.g., addition of floating point types) on the same GPU device. | |
| * However, results for pseudo-associative reduction may be inconsistent | |
| * from one device to a another device of a different compute-capability | |
| * because CUB can employ different tile-sizing for different architectures. | |
| * - \devicestorage | |
| * | |
| * \par Snippet | |
| * The code snippet below illustrates the argmin-reduction of a device vector of \p int data elements. | |
| * \par | |
| * \code | |
| * #include <cub/cub.cuh> // or equivalently <cub/device/device_radix_sort.cuh> | |
| * | |
| * // Declare, allocate, and initialize device-accessible pointers for input and output | |
| * int num_items; // e.g., 7 | |
| * int *d_in; // e.g., [8, 6, 7, 5, 3, 0, 9] | |
| * KeyValuePair<int, int> *d_out; // e.g., [{-,-}] | |
| * ... | |
| * | |
| * // Determine temporary device storage requirements | |
| * void *d_temp_storage = NULL; | |
| * size_t temp_storage_bytes = 0; | |
| * cub::DeviceReduce::ArgMin(d_temp_storage, temp_storage_bytes, d_in, d_argmin, num_items); | |
| * | |
| * // Allocate temporary storage | |
| * cudaMalloc(&d_temp_storage, temp_storage_bytes); | |
| * | |
| * // Run argmin-reduction | |
| * cub::DeviceReduce::ArgMin(d_temp_storage, temp_storage_bytes, d_in, d_argmin, num_items); | |
| * | |
| * // d_out <-- [{5, 0}] | |
| * | |
| * \endcode | |
| * | |
| * \tparam InputIteratorT <b>[inferred]</b> Random-access input iterator type for reading input items (of some type \p T) \iterator | |
| * \tparam OutputIteratorT <b>[inferred]</b> Output iterator type for recording the reduced aggregate (having value type <tt>cub::KeyValuePair<int, T></tt>) \iterator | |
| */ | |
| template < | |
| typename InputIteratorT, | |
| typename OutputIteratorT> | |
| CUB_RUNTIME_FUNCTION | |
| static cudaError_t ArgMin( | |
| void *d_temp_storage, ///< [in] %Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done. | |
| size_t &temp_storage_bytes, ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation | |
| InputIteratorT d_in, ///< [in] Pointer to the input sequence of data items | |
| OutputIteratorT d_out, ///< [out] Pointer to the output aggregate | |
| int num_items, ///< [in] Total number of input items (i.e., length of \p d_in) | |
| cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>. | |
| bool debug_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Also causes launch configurations to be printed to the console. Default is \p 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); | |
| } | |
| /** | |
| * \brief Computes a device-wide maximum using the greater-than ('>') operator. | |
| * | |
| * \par | |
| * - Uses <tt>std::numeric_limits<T>::lowest()</tt> as the initial value of the reduction. | |
| * - Does not support \p > operators that are non-commutative. | |
| * - Provides "run-to-run" determinism for pseudo-associative reduction | |
| * (e.g., addition of floating point types) on the same GPU device. | |
| * However, results for pseudo-associative reduction may be inconsistent | |
| * from one device to a another device of a different compute-capability | |
| * because CUB can employ different tile-sizing for different architectures. | |
| * - \devicestorage | |
| * | |
| * \par Snippet | |
| * The code snippet below illustrates the max-reduction of a device vector of \p int data elements. | |
| * \par | |
| * \code | |
| * #include <cub/cub.cuh> // or equivalently <cub/device/device_radix_sort.cuh> | |
| * | |
| * // Declare, allocate, and initialize device-accessible pointers for input and output | |
| * int num_items; // e.g., 7 | |
| * int *d_in; // e.g., [8, 6, 7, 5, 3, 0, 9] | |
| * int *d_out; // e.g., [-] | |
| * ... | |
| * | |
| * // Determine temporary device storage requirements | |
| * void *d_temp_storage = NULL; | |
| * size_t temp_storage_bytes = 0; | |
| * cub::DeviceReduce::Max(d_temp_storage, temp_storage_bytes, d_in, d_max, num_items); | |
| * | |
| * // Allocate temporary storage | |
| * cudaMalloc(&d_temp_storage, temp_storage_bytes); | |
| * | |
| * // Run max-reduction | |
| * cub::DeviceReduce::Max(d_temp_storage, temp_storage_bytes, d_in, d_max, num_items); | |
| * | |
| * // d_out <-- [9] | |
| * | |
| * \endcode | |
| * | |
| * \tparam InputIteratorT <b>[inferred]</b> Random-access input iterator type for reading input items \iterator | |
| * \tparam OutputIteratorT <b>[inferred]</b> Output iterator type for recording the reduced aggregate \iterator | |
| */ | |
| template < | |
| typename InputIteratorT, | |
| typename OutputIteratorT> | |
| CUB_RUNTIME_FUNCTION | |
| static cudaError_t Max( | |
| void *d_temp_storage, ///< [in] %Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done. | |
| size_t &temp_storage_bytes, ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation | |
| InputIteratorT d_in, ///< [in] Pointer to the input sequence of data items | |
| OutputIteratorT d_out, ///< [out] Pointer to the output aggregate | |
| int num_items, ///< [in] Total number of input items (i.e., length of \p d_in) | |
| cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>. | |
| bool debug_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Also causes launch configurations to be printed to the console. Default is \p 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); | |
| } | |
| /** | |
| * \brief Finds the first device-wide maximum using the greater-than ('>') operator, also returning the index of that item | |
| * | |
| * \par | |
| * - The output value type of \p d_out is cub::KeyValuePair <tt><int, T></tt> (assuming the value type of \p d_in is \p T) | |
| * - The maximum is written to <tt>d_out.value</tt> and its offset in the input array is written to <tt>d_out.key</tt>. | |
| * - The <tt>{1, std::numeric_limits<T>::lowest()}</tt> tuple is produced for zero-length inputs | |
| * - Does not support \p > operators that are non-commutative. | |
| * - Provides "run-to-run" determinism for pseudo-associative reduction | |
| * (e.g., addition of floating point types) on the same GPU device. | |
| * However, results for pseudo-associative reduction may be inconsistent | |
| * from one device to a another device of a different compute-capability | |
| * because CUB can employ different tile-sizing for different architectures. | |
| * - \devicestorage | |
| * | |
| * \par Snippet | |
| * The code snippet below illustrates the argmax-reduction of a device vector of \p int data elements. | |
| * \par | |
| * \code | |
| * #include <cub/cub.cuh> // or equivalently <cub/device/device_reduce.cuh> | |
| * | |
| * // Declare, allocate, and initialize device-accessible pointers for input and output | |
| * int num_items; // e.g., 7 | |
| * int *d_in; // e.g., [8, 6, 7, 5, 3, 0, 9] | |
| * KeyValuePair<int, int> *d_out; // e.g., [{-,-}] | |
| * ... | |
| * | |
| * // Determine temporary device storage requirements | |
| * void *d_temp_storage = NULL; | |
| * size_t temp_storage_bytes = 0; | |
| * cub::DeviceReduce::ArgMax(d_temp_storage, temp_storage_bytes, d_in, d_argmax, num_items); | |
| * | |
| * // Allocate temporary storage | |
| * cudaMalloc(&d_temp_storage, temp_storage_bytes); | |
| * | |
| * // Run argmax-reduction | |
| * cub::DeviceReduce::ArgMax(d_temp_storage, temp_storage_bytes, d_in, d_argmax, num_items); | |
| * | |
| * // d_out <-- [{6, 9}] | |
| * | |
| * \endcode | |
| * | |
| * \tparam InputIteratorT <b>[inferred]</b> Random-access input iterator type for reading input items (of some type \p T) \iterator | |
| * \tparam OutputIteratorT <b>[inferred]</b> Output iterator type for recording the reduced aggregate (having value type <tt>cub::KeyValuePair<int, T></tt>) \iterator | |
| */ | |
| template < | |
| typename InputIteratorT, | |
| typename OutputIteratorT> | |
| CUB_RUNTIME_FUNCTION | |
| static cudaError_t ArgMax( | |
| void *d_temp_storage, ///< [in] %Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done. | |
| size_t &temp_storage_bytes, ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation | |
| InputIteratorT d_in, ///< [in] Pointer to the input sequence of data items | |
| OutputIteratorT d_out, ///< [out] Pointer to the output aggregate | |
| int num_items, ///< [in] Total number of input items (i.e., length of \p d_in) | |
| cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>. | |
| bool debug_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Also causes launch configurations to be printed to the console. Default is \p 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); | |
| } | |
| /** | |
| * \brief Reduces segments of values, where segments are demarcated by corresponding runs of identical keys. | |
| * | |
| * \par | |
| * This operation computes segmented reductions within \p d_values_in using | |
| * the specified binary \p reduction_op functor. The segments are identified by | |
| * "runs" of corresponding keys in \p d_keys_in, where runs are maximal ranges of | |
| * consecutive, identical keys. For the <em>i</em><sup>th</sup> run encountered, | |
| * the first key of the run and the corresponding value aggregate of that run are | |
| * written to <tt>d_unique_out[<em>i</em>]</tt> and <tt>d_aggregates_out[<em>i</em>]</tt>, | |
| * respectively. The total number of runs encountered is written to \p d_num_runs_out. | |
| * | |
| * \par | |
| * - The <tt>==</tt> equality operator is used to determine whether keys are equivalent | |
| * - Provides "run-to-run" determinism for pseudo-associative reduction | |
| * (e.g., addition of floating point types) on the same GPU device. | |
| * However, results for pseudo-associative reduction may be inconsistent | |
| * from one device to a another device of a different compute-capability | |
| * because CUB can employ different tile-sizing for different architectures. | |
| * - \devicestorage | |
| * | |
| * \par Performance | |
| * The following chart illustrates reduction-by-key (sum) performance across | |
| * different CUDA architectures for \p fp32 and \p fp64 values, respectively. Segments | |
| * are identified by \p int32 keys, and have lengths uniformly sampled from [1,1000]. | |
| * | |
| * \image html reduce_by_key_fp32_len_500.png | |
| * \image html reduce_by_key_fp64_len_500.png | |
| * | |
| * \par | |
| * The following charts are similar, but with segment lengths uniformly sampled from [1,10]: | |
| * | |
| * \image html reduce_by_key_fp32_len_5.png | |
| * \image html reduce_by_key_fp64_len_5.png | |
| * | |
| * \par Snippet | |
| * The code snippet below illustrates the segmented reduction of \p int values grouped | |
| * by runs of associated \p int keys. | |
| * \par | |
| * \code | |
| * #include <cub/cub.cuh> // or equivalently <cub/device/device_reduce.cuh> | |
| * | |
| * // CustomMin functor | |
| * struct CustomMin | |
| * { | |
| * template <typename T> | |
| * CUB_RUNTIME_FUNCTION __forceinline__ | |
| * T operator()(const T &a, const T &b) const { | |
| * return (b < a) ? b : a; | |
| * } | |
| * }; | |
| * | |
| * // Declare, allocate, and initialize device-accessible pointers for input and output | |
| * int num_items; // e.g., 8 | |
| * int *d_keys_in; // e.g., [0, 2, 2, 9, 5, 5, 5, 8] | |
| * int *d_values_in; // e.g., [0, 7, 1, 6, 2, 5, 3, 4] | |
| * int *d_unique_out; // e.g., [-, -, -, -, -, -, -, -] | |
| * int *d_aggregates_out; // e.g., [-, -, -, -, -, -, -, -] | |
| * int *d_num_runs_out; // e.g., [-] | |
| * CustomMin reduction_op; | |
| * ... | |
| * | |
| * // Determine temporary device storage requirements | |
| * void *d_temp_storage = NULL; | |
| * size_t temp_storage_bytes = 0; | |
| * cub::DeviceReduce::ReduceByKey(d_temp_storage, temp_storage_bytes, d_keys_in, d_unique_out, d_values_in, d_aggregates_out, d_num_runs_out, reduction_op, num_items); | |
| * | |
| * // Allocate temporary storage | |
| * cudaMalloc(&d_temp_storage, temp_storage_bytes); | |
| * | |
| * // Run reduce-by-key | |
| * cub::DeviceReduce::ReduceByKey(d_temp_storage, temp_storage_bytes, d_keys_in, d_unique_out, d_values_in, d_aggregates_out, d_num_runs_out, reduction_op, num_items); | |
| * | |
| * // d_unique_out <-- [0, 2, 9, 5, 8] | |
| * // d_aggregates_out <-- [0, 1, 6, 2, 4] | |
| * // d_num_runs_out <-- [5] | |
| * | |
| * \endcode | |
| * | |
| * \tparam KeysInputIteratorT <b>[inferred]</b> Random-access input iterator type for reading input keys \iterator | |
| * \tparam UniqueOutputIteratorT <b>[inferred]</b> Random-access output iterator type for writing unique output keys \iterator | |
| * \tparam ValuesInputIteratorT <b>[inferred]</b> Random-access input iterator type for reading input values \iterator | |
| * \tparam AggregatesOutputIterator <b>[inferred]</b> Random-access output iterator type for writing output value aggregates \iterator | |
| * \tparam NumRunsOutputIteratorT <b>[inferred]</b> Output iterator type for recording the number of runs encountered \iterator | |
| * \tparam ReductionOpT <b>[inferred]</b> Binary reduction functor type having member <tt>T operator()(const T &a, const T &b)</tt> | |
| */ | |
| 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, ///< [in] %Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done. | |
| size_t &temp_storage_bytes, ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation | |
| KeysInputIteratorT d_keys_in, ///< [in] Pointer to the input sequence of keys | |
| UniqueOutputIteratorT d_unique_out, ///< [out] Pointer to the output sequence of unique keys (one key per run) | |
| ValuesInputIteratorT d_values_in, ///< [in] Pointer to the input sequence of corresponding values | |
| AggregatesOutputIteratorT d_aggregates_out, ///< [out] Pointer to the output sequence of value aggregates (one aggregate per run) | |
| NumRunsOutputIteratorT d_num_runs_out, ///< [out] Pointer to total number of runs encountered (i.e., the length of d_unique_out) | |
| ReductionOpT reduction_op, ///< [in] Binary reduction functor | |
| int num_items, ///< [in] Total number of associated key+value pairs (i.e., the length of \p d_in_keys and \p d_in_values) | |
| cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>. | |
| bool debug_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is \p 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); | |
| } | |
| }; | |
| /** | |
| * \example example_device_reduce.cu | |
| */ | |
| } // CUB namespace | |
| CUB_NS_POSTFIX // Optional outer namespace(s) | |