Alien-XGBoost
view release on metacpan or search on metacpan
xgboost/cub/cub/device/device_segmented_reduce.cuh view on Meta::CPAN
* #include <cub/cub.cuh> // or equivalently <cub/device/device_radix_sort.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_segments; // e.g., 3
* int *d_offsets; // e.g., [0, 3, 3, 7]
* int *d_in; // e.g., [8, 6, 7, 5, 3, 0, 9]
* int *d_out; // e.g., [-, -, -]
* CustomMin min_op;
* int initial_value; // e.g., INT_MAX
* ...
*
* // Determine temporary device storage requirements
* void *d_temp_storage = NULL;
* size_t temp_storage_bytes = 0;
* cub::DeviceSegmentedReduce::Reduce(d_temp_storage, temp_storage_bytes, d_in, d_out,
* num_segments, d_offsets, d_offsets + 1, min_op, initial_value);
*
* // Allocate temporary storage
* cudaMalloc(&d_temp_storage, temp_storage_bytes);
*
* // Run reduction
* cub::DeviceSegmentedReduce::Reduce(d_temp_storage, temp_storage_bytes, d_in, d_out,
* num_segments, d_offsets, d_offsets + 1, min_op, initial_value);
*
* // d_out <-- [6, INT_MAX, 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 ReductionOp <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 ReductionOp,
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_segments, ///< [in] The number of segments that comprise the sorting data
int *d_begin_offsets, ///< [in] %Device-accessible pointer to the sequence of beginning offsets of length \p num_segments, such that <tt>d_begin_offsets[i]</tt> is the first element of the <em>i</em><sup>th</...
int *d_end_offsets, ///< [in] %Device-accessible pointer to the sequence of ending offsets of length \p num_segments, such that <tt>d_end_offsets[i]-1</tt> is the last element of the <em>i</em><sup>th</sup>...
ReductionOp reduction_op, ///< [in] Binary reduction functor
T initial_value, ///< [in] Initial value of the reduction for each segment
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. Defa...
{
// Signed integer type for global offsets
typedef int OffsetT;
return DispatchSegmentedReduce<InputIteratorT, OutputIteratorT, 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);
}
/**
* \brief Computes a device-wide segmented sum using the addition ('+') operator.
*
* \par
* - Uses \p 0 as the initial value of the reduction for each segment.
* - When input a contiguous sequence of segments, a single sequence
* \p segment_offsets (of length <tt>num_segments+1</tt>) can be aliased
* for both the \p d_begin_offsets and \p d_end_offsets parameters (where
* the latter is specified as <tt>segment_offsets+1</tt>).
* - Does not support \p + operators that are non-commutative..
* - \devicestorage
*
* \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_segments; // e.g., 3
* int *d_offsets; // e.g., [0, 3, 3, 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::DeviceSegmentedReduce::Sum(d_temp_storage, temp_storage_bytes, d_in, d_out,
* num_segments, d_offsets, d_offsets + 1);
*
* // Allocate temporary storage
* cudaMalloc(&d_temp_storage, temp_storage_bytes);
*
* // Run sum-reduction
* cub::DeviceSegmentedReduce::Sum(d_temp_storage, temp_storage_bytes, d_in, d_out,
* num_segments, d_offsets, d_offsets + 1);
*
* // d_out <-- [21, 0, 17]
*
* \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_segments, ///< [in] The number of segments that comprise the sorting data
int *d_begin_offsets, ///< [in] %Device-accessible pointer to the sequence of beginning offsets of length \p num_segments, such that <tt>d_begin_offsets[i]</tt> is the first element of the <em>i</em><sup>th</...
int *d_end_offsets, ///< [in] %Device-accessible pointer to the sequence of ending offsets of length \p num_segments, such that <tt>d_end_offsets[i]-1</tt> is the last element of the <em>i</em><sup>th</sup>...
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. Defa...
{
// 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, 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);
}
/**
* \brief Computes a device-wide segmented minimum using the less-than ('<') operator.
*
* \par
* - Uses <tt>std::numeric_limits<T>::max()</tt> as the initial value of the reduction for each segment.
* - When input a contiguous sequence of segments, a single sequence
* \p segment_offsets (of length <tt>num_segments+1</tt>) can be aliased
* for both the \p d_begin_offsets and \p d_end_offsets parameters (where
* the latter is specified as <tt>segment_offsets+1</tt>).
* - Does not support \p < operators that are non-commutative.
* - \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_segments; // e.g., 3
* int *d_offsets; // e.g., [0, 3, 3, 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::DeviceSegmentedReduce::Min(d_temp_storage, temp_storage_bytes, d_in, d_out,
* num_segments, d_offsets, d_offsets + 1);
*
* // Allocate temporary storage
* cudaMalloc(&d_temp_storage, temp_storage_bytes);
*
* // Run min-reduction
* cub::DeviceSegmentedReduce::Min(d_temp_storage, temp_storage_bytes, d_in, d_out,
* num_segments, d_offsets, d_offsets + 1);
*
* // d_out <-- [6, INT_MAX, 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_segments, ///< [in] The number of segments that comprise the sorting data
int *d_begin_offsets, ///< [in] %Device-accessible pointer to the sequence of beginning offsets of length \p num_segments, such that <tt>d_begin_offsets[i]</tt> is the first element of the <em>i</em><sup>th</...
int *d_end_offsets, ///< [in] %Device-accessible pointer to the sequence of ending offsets of length \p num_segments, such that <tt>d_end_offsets[i]-1</tt> is the last element of the <em>i</em><sup>th</sup>...
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. Defa...
{
// 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, 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);
}
/**
* \brief Finds the first device-wide minimum in each segment using the less-than ('<') operator, also returning the in-segment 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 of the <em>i</em><sup>th</sup> segment is written to <tt>d_out[i].value</tt> and its offset in that segment is written to <tt>d_out[i].key</tt>.
* - The <tt>{1, std::numeric_limits<T>::max()}</tt> tuple is produced for zero-length inputs
* - When input a contiguous sequence of segments, a single sequence
* \p segment_offsets (of length <tt>num_segments+1</tt>) can be aliased
* for both the \p d_begin_offsets and \p d_end_offsets parameters (where
* the latter is specified as <tt>segment_offsets+1</tt>).
* - Does not support \p < operators that are non-commutative.
* - \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_segments; // e.g., 3
* int *d_offsets; // e.g., [0, 3, 3, 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::DeviceSegmentedReduce::ArgMin(d_temp_storage, temp_storage_bytes, d_in, d_out,
* num_segments, d_offsets, d_offsets + 1);
*
* // Allocate temporary storage
* cudaMalloc(&d_temp_storage, temp_storage_bytes);
*
* // Run argmin-reduction
* cub::DeviceSegmentedReduce::ArgMin(d_temp_storage, temp_storage_bytes, d_in, d_out,
* num_segments, d_offsets, d_offsets + 1);
*
* // d_out <-- [{1,6}, {1,INT_MAX}, {2,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>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_segments, ///< [in] The number of segments that comprise the sorting data
int *d_begin_offsets, ///< [in] %Device-accessible pointer to the sequence of beginning offsets of length \p num_segments, such that <tt>d_begin_offsets[i]</tt> is the first element of the <em>i</em><sup>th</...
int *d_end_offsets, ///< [in] %Device-accessible pointer to the sequence of ending offsets of length \p num_segments, such that <tt>d_end_offsets[i]-1</tt> is the last element of the <em>i</em><sup>th</sup>...
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. Defa...
{
// 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, 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);
}
/**
* \brief Computes a device-wide segmented maximum using the greater-than ('>') operator.
*
* \par
* - Uses <tt>std::numeric_limits<T>::lowest()</tt> as the initial value of the reduction.
* - When input a contiguous sequence of segments, a single sequence
* \p segment_offsets (of length <tt>num_segments+1</tt>) can be aliased
* for both the \p d_begin_offsets and \p d_end_offsets parameters (where
* the latter is specified as <tt>segment_offsets+1</tt>).
* - Does not support \p > operators that are non-commutative.
* - \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_segments; // e.g., 3
* int *d_offsets; // e.g., [0, 3, 3, 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::DeviceSegmentedReduce::Max(d_temp_storage, temp_storage_bytes, d_in, d_out,
* num_segments, d_offsets, d_offsets + 1);
*
* // Allocate temporary storage
* cudaMalloc(&d_temp_storage, temp_storage_bytes);
*
* // Run max-reduction
* cub::DeviceSegmentedReduce::Max(d_temp_storage, temp_storage_bytes, d_in, d_out,
* num_segments, d_offsets, d_offsets + 1);
*
* // d_out <-- [8, INT_MIN, 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_segments, ///< [in] The number of segments that comprise the sorting data
int *d_begin_offsets, ///< [in] %Device-accessible pointer to the sequence of beginning offsets of length \p num_segments, such that <tt>d_begin_offsets[i]</tt> is the first element of the <em>i</em><sup>th</...
int *d_end_offsets, ///< [in] %Device-accessible pointer to the sequence of ending offsets of length \p num_segments, such that <tt>d_end_offsets[i]-1</tt> is the last element of the <em>i</em><sup>th</sup>...
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. Defa...
{
// 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, 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);
}
/**
* \brief Finds the first device-wide maximum in each segment using the greater-than ('>') operator, also returning the in-segment 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 of the <em>i</em><sup>th</sup> segment is written to <tt>d_out[i].value</tt> and its offset in that segment is written to <tt>d_out[i].key</tt>.
* - The <tt>{1, std::numeric_limits<T>::lowest()}</tt> tuple is produced for zero-length inputs
* - When input a contiguous sequence of segments, a single sequence
* \p segment_offsets (of length <tt>num_segments+1</tt>) can be aliased
* for both the \p d_begin_offsets and \p d_end_offsets parameters (where
* the latter is specified as <tt>segment_offsets+1</tt>).
* - Does not support \p > operators that are non-commutative.
* - \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_segments; // e.g., 3
* int *d_offsets; // e.g., [0, 3, 3, 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::DeviceSegmentedReduce::ArgMax(d_temp_storage, temp_storage_bytes, d_in, d_out,
* num_segments, d_offsets, d_offsets + 1);
*
* // Allocate temporary storage
* cudaMalloc(&d_temp_storage, temp_storage_bytes);
*
* // Run argmax-reduction
* cub::DeviceSegmentedReduce::ArgMax(d_temp_storage, temp_storage_bytes, d_in, d_out,
* num_segments, d_offsets, d_offsets + 1);
*
* // d_out <-- [{0,8}, {1,INT_MIN}, {3,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>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_segments, ///< [in] The number of segments that comprise the sorting data
int *d_begin_offsets, ///< [in] %Device-accessible pointer to the sequence of beginning offsets of length \p num_segments, such that <tt>d_begin_offsets[i]</tt> is the first element of the <em>i</em><sup>th</...
int *d_end_offsets, ///< [in] %Device-accessible pointer to the sequence of ending offsets of length \p num_segments, such that <tt>d_end_offsets[i]-1</tt> is the last element of the <em>i</em><sup>th</sup>...
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. Defa...
{
// 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, 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)
( run in 1.368 second using v1.01-cache-2.11-cpan-39bf76dae61 )