<PRE>fix ID group-ID style bodystyle args keyword values ...
</PRE>
<UL><LI>ID, group-ID are documented in <A HREF = "fix.html">fix</A> command
<LI>style = <I>rigid</I> or <I>rigid/nve</I> or <I>rigid/nvt</I> or <I>rigid/npt</I> or <I>rigid/nph</I> or <I>rigid/small</I> or <I>rigid/nve/small</I> or <I>rigid/nvt/small</I> or <I>rigid/npt/small</I> or <I>rigid/nph/small</I>
<LI>bodystyle = <I>single</I> or <I>molecule</I> or <I>group</I>
<PRE> <I>single</I> args = none
<I>molecule</I> args = none
<I>group</I> args = N groupID1 groupID2 ...
N = # of groups
groupID1, groupID2, ... = list of N group IDs
</PRE>
<LI>zero or more keyword/value pairs may be appended
<LI>keyword = <I>langevin</I> or <I>temp</I> or <I>iso</I> or <I>aniso</I> or <I>x</I> or <I>y</I> or <I>z</I> or <I>couple</I> or <I>tparam</I> or <I>pchain</I> or <I>dilate</I> or <I>force</I> or <I>torque</I> or <I>infile</I>
fix ID group-ID style bodystyle args keyword values ... :pre
ID, group-ID are documented in "fix"_fix.html command :ulb,l
style = {rigid} or {rigid/nve} or {rigid/nvt} or {rigid/npt} or {rigid/nph} or {rigid/small} or {rigid/nve/small} or {rigid/nvt/small} or {rigid/npt/small} or {rigid/nph/small} :l
bodystyle = {single} or {molecule} or {group} :l
{single} args = none
{molecule} args = none
{group} args = N groupID1 groupID2 ...
N = # of groups
groupID1, groupID2, ... = list of N group IDs :pre
zero or more keyword/value pairs may be appended :l
keyword = {langevin} or {temp} or {iso} or {aniso} or {x} or {y} or {z} or {couple} or {tparam} or {pchain} or {dilate} or {force} or {torque} or {infile} :l
{langevin} values = Tstart Tstop Tperiod seed
Tstart,Tstop = desired temperature at start/stop of run (temperature units)
Tdamp = temperature damping parameter (time units)
seed = random number seed to use for white noise (positive integer)
{temp} values = Tstart Tstop Tdamp
Tstart,Tstop = desired temperature at start/stop of run (temperature units)
Tdamp = temperature damping parameter (time units)
{iso} or {aniso} values = Pstart Pstop Pdamp
Pstart,Pstop = scalar external pressure at start/end of run (pressure units)
Pdamp = pressure damping parameter (time units)
{x} or {y} or {z} values = Pstart Pstop Pdamp
Pstart,Pstop = external stress tensor component at start/end of run (pressure units)
Pdamp = stress damping parameter (time units)
{couple} = {none} or {xyz} or {xy} or {yz} or {xz}
{tparam} values = Tchain Titer Torder
Tchain = length of Nose/Hoover thermostat chain
Titer = number of thermostat iterations performed
Torder = 3 or 5 = Yoshida-Suzuki integration parameters
{pchain} values = Pchain
Pchain = length of the Nose/Hoover thermostat chain coupled with the barostat
{dilate} value = dilate-group-ID
dilate-group-ID = only dilate atoms in this group due to barostat volume changes
{force} values = M xflag yflag zflag
M = which rigid body from 1-Nbody (see asterisk form below)
xflag,yflag,zflag = off/on if component of center-of-mass force is active
{torque} values = M xflag yflag zflag
M = which rigid body from 1-Nbody (see asterisk form below)
xflag,yflag,zflag = off/on if component of center-of-mass torque is active
{infile} filename
filename = file with per-body values of mass, center-of-mass, moments of inertia
{mol} value = template-ID
template-ID = ID of molecule template specified in a separate "molecule"_molecule.html command :pre
:ule
[Examples:]
fix 1 clump rigid single
fix 1 clump rigid/small molecule
fix 1 clump rigid single force 1 off off on langevin 1.0 1.0 1.0 428984
- * The cub::BlockDiscontinuity class provides [<em>collective</em>](index.html#sec0) methods for flagging discontinuities within an ordered set of items partitioned across a CUDA thread block.
- */
-
-#pragma once
-
-#include "../util_type.cuh"
-#include "../util_namespace.cuh"
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-/**
- * \brief The BlockDiscontinuity class provides [<em>collective</em>](index.html#sec0) methods for flagging discontinuities within an ordered set of items partitioned across a CUDA thread block. 
- * \ingroup BlockModule
- *
- * \par Overview
- * A set of "head flags" (or "tail flags") is often used to indicate corresponding items
- * that differ from their predecessors (or successors). For example, head flags are convenient
- * for demarcating disjoint data segments as part of a segmented scan or reduction.
- *
- * \tparam T The data type to be flagged.
- * \tparam BLOCK_THREADS The thread block size in threads.
- *
- * \par A Simple Example
- * \blockcollective{BlockDiscontinuity}
- * \par
- * The code snippet below illustrates the head flagging of 512 integer items that
- * are partitioned in a [<em>blocked arrangement</em>](index.html#sec5sec4) across 128 threads
- * where each thread owns 4 consecutive items.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockDiscontinuity for 128 threads on type int
- * \brief Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockDiscontinuity()
- :
- temp_storage(PrivateStorage()),
- linear_tid(threadIdx.x)
- {}
-
-
- /**
- * \brief Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockDiscontinuity(
- TempStorage &temp_storage) ///< [in] Reference to memory allocation having layout type TempStorage
- :
- temp_storage(temp_storage.Alias()),
- linear_tid(threadIdx.x)
- {}
-
-
- /**
- * \brief Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier
- */
- __device__ __forceinline__ BlockDiscontinuity(
- int linear_tid) ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- :
- temp_storage(PrivateStorage()),
- linear_tid(linear_tid)
- {}
-
-
- /**
- * \brief Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier.
- */
- __device__ __forceinline__ BlockDiscontinuity(
- TempStorage &temp_storage, ///< [in] Reference to memory allocation having layout type TempStorage
- int linear_tid) ///< [in] <b>[optional]</b> A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- * \brief Sets head flags indicating discontinuities between items partitioned across the thread block, for which the first item has no reference and is always flagged.
- *
- * The flag <tt>head_flags<sub><em>i</em></sub></tt> is set for item
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam FlagT <b>[inferred]</b> The flag type (must be an integer type)
- * \tparam FlagOp <b>[inferred]</b> Binary predicate functor type having member <tt>T operator()(const T &a, const T &b)</tt> or member <tt>T operator()(const T &a, const T &b, unsigned int b_index)</tt>, and returning \p true if a discontinuity exists between \p a and \p b, otherwise \p false. \p b_index is the rank of b in the aggregate tile of data.
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam FlagT <b>[inferred]</b> The flag type (must be an integer type)
- * \tparam FlagOp <b>[inferred]</b> Binary predicate functor type having member <tt>T operator()(const T &a, const T &b)</tt> or member <tt>T operator()(const T &a, const T &b, unsigned int b_index)</tt>, and returning \p true if a discontinuity exists between \p a and \p b, otherwise \p false. \p b_index is the rank of b in the aggregate tile of data.
- T (&input)[ITEMS_PER_THREAD], ///< [in] Calling thread's input items
- FlagOp flag_op, ///< [in] Binary boolean flag predicate
- T tile_predecessor_item) ///< [in] <b>[<em>thread</em><sub>0</sub> only]</b> Item with which to compare the first tile item (<tt>input<sub>0</sub></tt> from <em>thread</em><sub>0</sub>).
- * \brief Sets tail flags indicating discontinuities between items partitioned across the thread block, for which the last item has no reference and is always flagged.
- *
- * The flag <tt>tail_flags<sub><em>i</em></sub></tt> is set for item
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam FlagT <b>[inferred]</b> The flag type (must be an integer type)
- * \tparam FlagOp <b>[inferred]</b> Binary predicate functor type having member <tt>T operator()(const T &a, const T &b)</tt> or member <tt>T operator()(const T &a, const T &b, unsigned int b_index)</tt>, and returning \p true if a discontinuity exists between \p a and \p b, otherwise \p false. \p b_index is the rank of b in the aggregate tile of data.
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam FlagT <b>[inferred]</b> The flag type (must be an integer type)
- * \tparam FlagOp <b>[inferred]</b> Binary predicate functor type having member <tt>T operator()(const T &a, const T &b)</tt> or member <tt>T operator()(const T &a, const T &b, unsigned int b_index)</tt>, and returning \p true if a discontinuity exists between \p a and \p b, otherwise \p false. \p b_index is the rank of b in the aggregate tile of data.
- T (&input)[ITEMS_PER_THREAD], ///< [in] Calling thread's input items
- FlagOp flag_op, ///< [in] Binary boolean flag predicate
- T tile_successor_item) ///< [in] <b>[<em>thread</em><sub><tt>BLOCK_THREADS</tt>-1</sub> only]</b> Item with which to compare the last tile item (<tt>input</tt><sub><em>ITEMS_PER_THREAD</em>-1</sub> from <em>thread</em><sub><em>BLOCK_THREADS</em>-1</sub>).
- {
- // Share first item
- temp_storage[linear_tid] = input[0];
-
- __syncthreads();
-
- // Set flag for last item
- int successor_item = (linear_tid == BLOCK_THREADS - 1) ?
- * The cub::BlockExchange class provides [<em>collective</em>](index.html#sec0) methods for rearranging data partitioned across a CUDA thread block.
- */
-
-#pragma once
-
-#include "../util_arch.cuh"
-#include "../util_macro.cuh"
-#include "../util_type.cuh"
-#include "../util_namespace.cuh"
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-/**
- * \brief The BlockExchange class provides [<em>collective</em>](index.html#sec0) methods for rearranging data partitioned across a CUDA thread block. 
- * \ingroup BlockModule
- *
- * \par Overview
- * It is commonplace for blocks of threads to rearrange data items between
- * threads. For example, the global memory subsystem prefers access patterns
- * where data items are "striped" across threads (where consecutive threads access consecutive items),
- * yet most block-wide operations prefer a "blocked" partitioning of items across threads
- * (where consecutive items belong to a single thread).
- *
- * \par
- * BlockExchange supports the following types of data exchanges:
- * - Transposing between [<em>blocked</em>](index.html#sec5sec4) and [<em>striped</em>](index.html#sec5sec4) arrangements
- * - Transposing between [<em>blocked</em>](index.html#sec5sec4) and [<em>warp-striped</em>](index.html#sec5sec4) arrangements
- * - Scattering ranked items to a [<em>blocked arrangement</em>](index.html#sec5sec4)
- * - Scattering ranked items to a [<em>striped arrangement</em>](index.html#sec5sec4)
- *
- * \tparam T The data type to be exchanged.
- * \tparam BLOCK_THREADS The thread block size in threads.
- * \tparam ITEMS_PER_THREAD The number of items partitioned onto each thread.
- * \tparam WARP_TIME_SLICING <b>[optional]</b> When \p true, only use enough shared memory for a single warp's worth of tile data, time-slicing the block-wide exchange over multiple synchronized rounds. Yields a smaller memory footprint at the expense of decreased parallelism. (Default: false)
- *
- * \par A Simple Example
- * \blockcollective{BlockExchange}
- * \par
- * The code snippet below illustrates the conversion from a "blocked" to a "striped" arrangement
- * of 512 integer items partitioned across 128 threads where each thread owns 4 items.
- * \brief Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockExchange()
- :
- temp_storage(PrivateStorage()),
- linear_tid(threadIdx.x),
- warp_lane(linear_tid & (WARP_THREADS - 1)),
- warp_id(linear_tid >> LOG_WARP_THREADS),
- warp_offset(warp_id * WARP_TIME_SLICED_ITEMS)
- {}
-
-
- /**
- * \brief Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockExchange(
- TempStorage &temp_storage) ///< [in] Reference to memory allocation having layout type TempStorage
- :
- temp_storage(temp_storage.Alias()),
- linear_tid(threadIdx.x),
- warp_lane(linear_tid & (WARP_THREADS - 1)),
- warp_id(linear_tid >> LOG_WARP_THREADS),
- warp_offset(warp_id * WARP_TIME_SLICED_ITEMS)
- {}
-
-
- /**
- * \brief Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier
- */
- __device__ __forceinline__ BlockExchange(
- int linear_tid) ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- :
- temp_storage(PrivateStorage()),
- linear_tid(linear_tid),
- warp_lane(linear_tid & (WARP_THREADS - 1)),
- warp_id(linear_tid >> LOG_WARP_THREADS),
- warp_offset(warp_id * WARP_TIME_SLICED_ITEMS)
- {}
-
-
- /**
- * \brief Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier.
- */
- __device__ __forceinline__ BlockExchange(
- TempStorage &temp_storage, ///< [in] Reference to memory allocation having layout type TempStorage
- int linear_tid) ///< [in] <b>[optional]</b> A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- * The cub::BlockHistogram class provides [<em>collective</em>](index.html#sec0) methods for constructing block-wide histograms from data samples partitioned across a CUDA thread block.
- * \brief The BlockHistogram class provides [<em>collective</em>](index.html#sec0) methods for constructing block-wide histograms from data samples partitioned across a CUDA thread block. 
- * \ingroup BlockModule
- *
- * \par Overview
- * A <a href="http://en.wikipedia.org/wiki/Histogram"><em>histogram</em></a>
- * counts the number of observations that fall into each of the disjoint categories (known as <em>bins</em>).
- *
- * \par
- * Optionally, BlockHistogram can be specialized to use different algorithms:
- * -# <b>cub::BLOCK_HISTO_SORT</b>. Sorting followed by differentiation. [More...](\ref cub::BlockHistogramAlgorithm)
- * -# <b>cub::BLOCK_HISTO_ATOMIC</b>. Use atomic addition to update byte counts directly. [More...](\ref cub::BlockHistogramAlgorithm)
- *
- * \tparam T The sample type being histogrammed (must be castable to an integer bin identifier)
- * \tparam BLOCK_THREADS The thread block size in threads
- * \tparam ITEMS_PER_THREAD The number of items per thread
- * \tparam BINS The number bins within the histogram
- * \tparam ALGORITHM <b>[optional]</b> cub::BlockHistogramAlgorithm enumerator specifying the underlying algorithm to use (default: cub::BLOCK_HISTO_SORT)
- *
- * \par A Simple Example
- * \blockcollective{BlockHistogram}
- * \par
- * The code snippet below illustrates a 256-bin histogram of 512 integer samples that
- * are partitioned across 128 threads where each thread owns 4 samples.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize a 256-bin BlockHistogram type for 128 threads having 4 character samples each
- * \brief Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockHistogram()
- :
- temp_storage(PrivateStorage()),
- linear_tid(threadIdx.x)
- {}
-
-
- /**
- * \brief Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockHistogram(
- TempStorage &temp_storage) ///< [in] Reference to memory allocation having layout type TempStorage
- :
- temp_storage(temp_storage.Alias()),
- linear_tid(threadIdx.x)
- {}
-
-
- /**
- * \brief Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier
- */
- __device__ __forceinline__ BlockHistogram(
- int linear_tid) ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- :
- temp_storage(PrivateStorage()),
- linear_tid(linear_tid)
- {}
-
-
- /**
- * \brief Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier.
- */
- __device__ __forceinline__ BlockHistogram(
- TempStorage &temp_storage, ///< [in] Reference to memory allocation having layout type TempStorage
- int linear_tid) ///< [in] <b>[optional]</b> A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- * \tparam T <b>[inferred]</b> The data type to load.
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam InputIteratorRA <b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type).
- */
-template <
- PtxLoadModifier MODIFIER,
- typename T,
- int ITEMS_PER_THREAD,
- typename InputIteratorRA>
-__device__ __forceinline__ void LoadBlocked(
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- InputIteratorRA block_itr, ///< [in] The thread block's base input iterator for loading from
- T (&items)[ITEMS_PER_THREAD]) ///< [out] Data to load
-{
- // Load directly in thread-blocked order
- #pragma unroll
- for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
- * \tparam T <b>[inferred]</b> The data type to load.
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam InputIteratorRA <b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type).
- */
-template <
- PtxLoadModifier MODIFIER,
- typename T,
- int ITEMS_PER_THREAD,
- typename InputIteratorRA>
-__device__ __forceinline__ void LoadBlocked(
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- InputIteratorRA block_itr, ///< [in] The thread block's base input iterator for loading from
- T (&items)[ITEMS_PER_THREAD], ///< [out] Data to load
- int valid_items) ///< [in] Number of valid items to load
-{
- int bounds = valid_items - (linear_tid * ITEMS_PER_THREAD);
-
- #pragma unroll
- for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
- * \brief Load a linear segment of items into a blocked arrangement across the thread block using the specified cache modifier, guarded by range, with a fall-back assignment of out-of-bound elements..
- * \tparam T <b>[inferred]</b> The data type to load.
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam InputIteratorRA <b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type).
- */
-template <
- PtxLoadModifier MODIFIER,
- typename T,
- int ITEMS_PER_THREAD,
- typename InputIteratorRA>
-__device__ __forceinline__ void LoadBlocked(
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- InputIteratorRA block_itr, ///< [in] The thread block's base input iterator for loading from
- T (&items)[ITEMS_PER_THREAD], ///< [out] Data to load
- int valid_items, ///< [in] Number of valid items to load
- T oob_default) ///< [in] Default value to assign out-of-bound items
-{
- int bounds = valid_items - (linear_tid * ITEMS_PER_THREAD);
-
- #pragma unroll
- for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
- * \tparam BLOCK_THREADS The thread block size in threads
- * \tparam T <b>[inferred]</b> The data type to load.
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam InputIteratorRA <b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type).
- */
-template <
- PtxLoadModifier MODIFIER,
- int BLOCK_THREADS,
- typename T,
- int ITEMS_PER_THREAD,
- typename InputIteratorRA>
-__device__ __forceinline__ void LoadStriped(
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- InputIteratorRA block_itr, ///< [in] The thread block's base input iterator for loading from
- T (&items)[ITEMS_PER_THREAD]) ///< [out] Data to load
-{
- #pragma unroll
- for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
- * \tparam BLOCK_THREADS The thread block size in threads
- * \tparam T <b>[inferred]</b> The data type to load.
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam InputIteratorRA <b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type).
- */
-template <
- PtxLoadModifier MODIFIER,
- int BLOCK_THREADS,
- typename T,
- int ITEMS_PER_THREAD,
- typename InputIteratorRA>
-__device__ __forceinline__ void LoadStriped(
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- InputIteratorRA block_itr, ///< [in] The thread block's base input iterator for loading from
- T (&items)[ITEMS_PER_THREAD], ///< [out] Data to load
- int valid_items) ///< [in] Number of valid items to load
-{
- int bounds = valid_items - linear_tid;
-
- #pragma unroll
- for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
- * \brief Load a linear segment of items into a striped arrangement across the thread block using the specified cache modifier, guarded by range, with a fall-back assignment of out-of-bound elements.
- * \tparam BLOCK_THREADS The thread block size in threads
- * \tparam T <b>[inferred]</b> The data type to load.
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam InputIteratorRA <b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type).
- */
-template <
- PtxLoadModifier MODIFIER,
- int BLOCK_THREADS,
- typename T,
- int ITEMS_PER_THREAD,
- typename InputIteratorRA>
-__device__ __forceinline__ void LoadStriped(
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- InputIteratorRA block_itr, ///< [in] The thread block's base input iterator for loading from
- T (&items)[ITEMS_PER_THREAD], ///< [out] Data to load
- int valid_items, ///< [in] Number of valid items to load
- T oob_default) ///< [in] Default value to assign out-of-bound items
-{
- int bounds = valid_items - linear_tid;
-
- #pragma unroll
- for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
- * \tparam T <b>[inferred]</b> The data type to load.
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam InputIteratorRA <b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type).
- */
-template <
- PtxLoadModifier MODIFIER,
- typename T,
- int ITEMS_PER_THREAD,
- typename InputIteratorRA>
-__device__ __forceinline__ void LoadWarpStriped(
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- InputIteratorRA block_itr, ///< [in] The thread block's base input iterator for loading from
- T (&items)[ITEMS_PER_THREAD]) ///< [out] Data to load
-{
- int tid = linear_tid & (PtxArchProps::WARP_THREADS - 1);
- int wid = linear_tid >> PtxArchProps::LOG_WARP_THREADS;
- int warp_offset = wid * PtxArchProps::WARP_THREADS * ITEMS_PER_THREAD;
-
- // Load directly in warp-striped order
- #pragma unroll
- for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
- * \tparam T <b>[inferred]</b> The data type to load.
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam InputIteratorRA <b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type).
- */
-template <
- PtxLoadModifier MODIFIER,
- typename T,
- int ITEMS_PER_THREAD,
- typename InputIteratorRA>
-__device__ __forceinline__ void LoadWarpStriped(
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- InputIteratorRA block_itr, ///< [in] The thread block's base input iterator for loading from
- T (&items)[ITEMS_PER_THREAD], ///< [out] Data to load
- int valid_items) ///< [in] Number of valid items to load
-{
- int tid = linear_tid & (PtxArchProps::WARP_THREADS - 1);
- int wid = linear_tid >> PtxArchProps::LOG_WARP_THREADS;
- int warp_offset = wid * PtxArchProps::WARP_THREADS * ITEMS_PER_THREAD;
- int bounds = valid_items - warp_offset - tid;
-
- // Load directly in warp-striped order
- #pragma unroll
- for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
- {
- if ((ITEM * PtxArchProps::WARP_THREADS) < bounds)
- * \brief Load a linear segment of items into a warp-striped arrangement across the thread block using the specified cache modifier, guarded by range, with a fall-back assignment of out-of-bound elements.
- *
- * \warpstriped
- *
- * \par Usage Considerations
- * The number of threads in the thread block must be a multiple of the architecture's warp size.
- * \tparam T <b>[inferred]</b> The data type to load.
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam InputIteratorRA <b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type).
- */
-template <
- PtxLoadModifier MODIFIER,
- typename T,
- int ITEMS_PER_THREAD,
- typename InputIteratorRA>
-__device__ __forceinline__ void LoadWarpStriped(
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- InputIteratorRA block_itr, ///< [in] The thread block's base input iterator for loading from
- T (&items)[ITEMS_PER_THREAD], ///< [out] Data to load
- int valid_items, ///< [in] Number of valid items to load
- T oob_default) ///< [in] Default value to assign out-of-bound items
-{
- int tid = linear_tid & (PtxArchProps::WARP_THREADS - 1);
- int wid = linear_tid >> PtxArchProps::LOG_WARP_THREADS;
- int warp_offset = wid * PtxArchProps::WARP_THREADS * ITEMS_PER_THREAD;
- int bounds = valid_items - warp_offset - tid;
-
- // Load directly in warp-striped order
- #pragma unroll
- for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- T *block_ptr, ///< [in] Input pointer for loading from
- T (&items)[ITEMS_PER_THREAD]) ///< [out] Data to load
-{
- enum
- {
- // Maximum CUDA vector size is 4 elements
- MAX_VEC_SIZE = CUB_MIN(4, ITEMS_PER_THREAD),
-
- // Vector size must be a power of two and an even divisor of the items per thread
- * \brief cub::BlockLoadAlgorithm enumerates alternative algorithms for cub::BlockLoad to read a linear segment of data from memory into a blocked arrangement across a CUDA thread block.
- */
-enum BlockLoadAlgorithm
-{
- /**
- * \par Overview
- *
- * A [<em>blocked arrangement</em>](index.html#sec5sec4) of data is read
- * directly from memory. The thread block reads items in a parallel "raking" fashion: thread<sub><em>i</em></sub>
- * reads the <em>i</em><sup>th</sup> segment of consecutive elements.
- *
- * \par Performance Considerations
- * - The utilization of memory transactions (coalescing) decreases as the
- * access stride between threads increases (i.e., the number items per thread).
- */
- BLOCK_LOAD_DIRECT,
-
- /**
- * \par Overview
- *
- * A [<em>blocked arrangement</em>](index.html#sec5sec4) of data is read directly
- * from memory using CUDA's built-in vectorized loads as a coalescing optimization.
- * The thread block reads items in a parallel "raking" fashion: thread<sub><em>i</em></sub> uses vector loads to
- * read the <em>i</em><sup>th</sup> segment of consecutive elements.
- *
- * For example, <tt>ld.global.v4.s32</tt> instructions will be generated when \p T = \p int and \p ITEMS_PER_THREAD > 4.
- *
- * \par Performance Considerations
- * - The utilization of memory transactions (coalescing) remains high until the the
- * access stride between threads (i.e., the number items per thread) exceeds the
- * maximum vector load width (typically 4 items or 64B, whichever is lower).
- * - The following conditions will prevent vectorization and loading will fall back to cub::BLOCK_LOAD_DIRECT:
- * - \p ITEMS_PER_THREAD is odd
- * - The \p InputIteratorRA is not a simple pointer type
- * - The block input offset is not quadword-aligned
- * - The data type \p T is not a built-in primitive or CUDA vector type (e.g., \p short, \p int2, \p double, \p float2, etc.)
- */
- BLOCK_LOAD_VECTORIZE,
-
- /**
- * \par Overview
- *
- * A [<em>striped arrangement</em>](index.html#sec5sec4) of data is read
- * directly from memory and then is locally transposed into a
- * [<em>blocked arrangement</em>](index.html#sec5sec4). The thread block
- * reads items in a parallel "strip-mining" fashion:
- * thread<sub><em>i</em></sub> reads items having stride \p BLOCK_THREADS
- * between them. cub::BlockExchange is then used to locally reorder the items
- * into a [<em>blocked arrangement</em>](index.html#sec5sec4).
- *
- * \par Performance Considerations
- * - The utilization of memory transactions (coalescing) remains high regardless
- * of items loaded per thread.
- * - The local reordering incurs slightly longer latencies and throughput than the
- * direct cub::BLOCK_LOAD_DIRECT and cub::BLOCK_LOAD_VECTORIZE alternatives.
- */
- BLOCK_LOAD_TRANSPOSE,
-
-
- /**
- * \par Overview
- *
- * A [<em>warp-striped arrangement</em>](index.html#sec5sec4) of data is read
- * directly from memory and then is locally transposed into a
- * [<em>blocked arrangement</em>](index.html#sec5sec4). Each warp reads its own
- * contiguous segment in a parallel "strip-mining" fashion: lane<sub><em>i</em></sub>
- * reads items having stride \p WARP_THREADS between them. cub::BlockExchange
- * is then used to locally reorder the items into a
- * - BLOCK_THREADS must be a multiple of WARP_THREADS
- *
- * \par Performance Considerations
- * - The utilization of memory transactions (coalescing) remains high regardless
- * of items loaded per thread.
- * - The local reordering incurs slightly longer latencies and throughput than the
- * direct cub::BLOCK_LOAD_DIRECT and cub::BLOCK_LOAD_VECTORIZE alternatives.
- */
- BLOCK_LOAD_WARP_TRANSPOSE,
-};
-
-
-/**
- * \brief The BlockLoad class provides [<em>collective</em>](index.html#sec0) data movement methods for loading a linear segment of items from memory into a [<em>blocked arrangement</em>](index.html#sec5sec4) across a CUDA thread block. 
- * \ingroup BlockModule
- *
- * \par Overview
- * The BlockLoad class provides a single data movement abstraction that can be specialized
- * to implement different cub::BlockLoadAlgorithm strategies. This facilitates different
- * performance policies for different architectures, data types, granularity sizes, etc.
- *
- * \par
- * Optionally, BlockLoad can be specialized by different data movement strategies:
- * -# <b>cub::BLOCK_LOAD_DIRECT</b>. A [<em>blocked arrangement</em>](index.html#sec5sec4)
- * of data is read directly from memory. [More...](\ref cub::BlockLoadAlgorithm)
- * -# <b>cub::BLOCK_LOAD_VECTORIZE</b>. A [<em>blocked arrangement</em>](index.html#sec5sec4)
- * of data is read directly from memory using CUDA's built-in vectorized loads as a
- * \tparam WARP_TIME_SLICING <b>[optional]</b> For transposition-based cub::BlockLoadAlgorithm parameterizations that utilize shared memory: When \p true, only use enough shared memory for a single warp's worth of data, time-slicing the block-wide exchange over multiple synchronized rounds (default: false)
- *
- * \par A Simple Example
- * \blockcollective{BlockLoad}
- * \par
- * The code snippet below illustrates the loading of a linear
- * segment of 512 integers into a "blocked" arrangement across 128 threads where each
- * thread owns 4 consecutive items. The load is specialized for \p BLOCK_LOAD_WARP_TRANSPOSE,
- * meaning memory references are efficiently coalesced using a warp-striped access
- * pattern (after which items are locally reordered among threads).
- * \brief Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockLoad()
- :
- temp_storage(PrivateStorage()),
- linear_tid(threadIdx.x)
- {}
-
-
- /**
- * \brief Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockLoad(
- TempStorage &temp_storage) ///< [in] Reference to memory allocation having layout type TempStorage
- :
- temp_storage(temp_storage.Alias()),
- linear_tid(threadIdx.x)
- {}
-
-
- /**
- * \brief Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier
- */
- __device__ __forceinline__ BlockLoad(
- int linear_tid) ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- :
- temp_storage(PrivateStorage()),
- linear_tid(linear_tid)
- {}
-
-
- /**
- * \brief Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier.
- */
- __device__ __forceinline__ BlockLoad(
- TempStorage &temp_storage, ///< [in] Reference to memory allocation having layout type TempStorage
- int linear_tid) ///< [in] <b>[optional]</b> A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- * cub::BlockRadixRank provides operations for ranking unsigned integer types within a CUDA threadblock
- */
-
-#pragma once
-
-#include "../util_arch.cuh"
-#include "../util_type.cuh"
-#include "../thread/thread_reduce.cuh"
-#include "../thread/thread_scan.cuh"
-#include "../block/block_scan.cuh"
-#include "../util_namespace.cuh"
-
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-/**
- * \brief BlockRadixRank provides operations for ranking unsigned integer types within a CUDA threadblock.
- * \ingroup BlockModule
- *
- * \par Overview
- * Blah...
- *
- * \tparam BLOCK_THREADS The thread block size in threads
- * \tparam RADIX_BITS <b>[optional]</b> The number of radix bits per digit place (default: 5 bits)
- * \tparam MEMOIZE_OUTER_SCAN <b>[optional]</b> Whether or not to buffer outer raking scan partials to incur fewer shared memory reads at the expense of higher register pressure (default: true for architectures SM35 and newer, false otherwise). See BlockScanAlgorithm::BLOCK_SCAN_RAKING_MEMOIZE for more details.
- * \tparam INNER_SCAN_ALGORITHM <b>[optional]</b> The cub::BlockScanAlgorithm algorithm to use (default: cub::BLOCK_SCAN_WARP_SCANS)
- * \brief Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockRadixRank()
- :
- temp_storage(PrivateStorage()),
- linear_tid(threadIdx.x)
- {}
-
-
- /**
- * \brief Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockRadixRank(
- TempStorage &temp_storage) ///< [in] Reference to memory allocation having layout type TempStorage
- :
- temp_storage(temp_storage.Alias()),
- linear_tid(threadIdx.x)
- {}
-
-
- /**
- * \brief Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier
- */
- __device__ __forceinline__ BlockRadixRank(
- int linear_tid) ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- :
- temp_storage(PrivateStorage()),
- linear_tid(linear_tid)
- {}
-
-
- /**
- * \brief Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier.
- */
- __device__ __forceinline__ BlockRadixRank(
- TempStorage &temp_storage, ///< [in] Reference to memory allocation having layout type TempStorage
- int linear_tid) ///< [in] <b>[optional]</b> A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- * The cub::BlockRadixSort class provides [<em>collective</em>](index.html#sec0) methods for radix sorting of items partitioned across a CUDA thread block.
- */
-
-
-#pragma once
-
-#include "../util_namespace.cuh"
-#include "../util_arch.cuh"
-#include "../util_type.cuh"
-#include "block_exchange.cuh"
-#include "block_radix_rank.cuh"
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-/**
- * \brief The cub::BlockRadixSort class provides [<em>collective</em>](index.html#sec0) methods for sorting items partitioned across a CUDA thread block using a radix sorting method. 
- * \ingroup BlockModule
- *
- * \par Overview
- * The [<em>radix sorting method</em>](http://en.wikipedia.org/wiki/Radix_sort) arranges
- * items into ascending order. It relies upon a positional representation for
- * keys, i.e., each key is comprised of an ordered sequence of symbols (e.g., digits,
- * characters, etc.) specified from least-significant to most-significant. For a
- * given input sequence of keys and a set of rules specifying a total ordering
- * of the symbolic alphabet, the radix sorting method produces a lexicographic
- * ordering of those keys.
- *
- * \par
- * BlockRadixSort can sort all of the built-in C++ numeric primitive types, e.g.:
- * <tt>unsigned char</tt>, \p int, \p double, etc. Within each key, the implementation treats fixed-length
- * bit-sequences of \p RADIX_BITS as radix digit places. Although the direct radix sorting
- * method can only be applied to unsigned integral types, BlockRadixSort
- * is able to sort signed and floating-point types via simple bit-wise transformations
- * that ensure lexicographic key ordering.
- *
- * \tparam Key Key type
- * \tparam BLOCK_THREADS The thread block size in threads
- * \tparam ITEMS_PER_THREAD The number of items per thread
- * \tparam Value <b>[optional]</b> Value type (default: cub::NullType)
- * \tparam RADIX_BITS <b>[optional]</b> The number of radix bits per digit place (default: 4 bits)
- * \tparam MEMOIZE_OUTER_SCAN <b>[optional]</b> Whether or not to buffer outer raking scan partials to incur fewer shared memory reads at the expense of higher register pressure (default: true for architectures SM35 and newer, false otherwise).
- * \tparam INNER_SCAN_ALGORITHM <b>[optional]</b> The cub::BlockScanAlgorithm algorithm to use (default: cub::BLOCK_SCAN_WARP_SCANS)
- * \brief Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockRadixSort()
- :
- temp_storage(PrivateStorage()),
- linear_tid(threadIdx.x)
- {}
-
-
- /**
- * \brief Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockRadixSort(
- TempStorage &temp_storage) ///< [in] Reference to memory allocation having layout type TempStorage
- :
- temp_storage(temp_storage.Alias()),
- linear_tid(threadIdx.x)
- {}
-
-
- /**
- * \brief Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier
- */
- __device__ __forceinline__ BlockRadixSort(
- int linear_tid) ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- :
- temp_storage(PrivateStorage()),
- linear_tid(linear_tid)
- {}
-
-
- /**
- * \brief Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier.
- */
- __device__ __forceinline__ BlockRadixSort(
- TempStorage &temp_storage, ///< [in] Reference to memory allocation having layout type TempStorage
- int linear_tid) ///< [in] <b>[optional]</b> A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- * \brief Performs a radix sort across a [<em>blocked arrangement</em>](index.html#sec5sec4) of keys, leaving them in a [<em>striped arrangement</em>](index.html#sec5sec4).
- *
- * \smemreuse
- *
- * The code snippet below illustrates a sort of 512 integer keys that
- * are initially partitioned in a [<em>blocked arrangement</em>](index.html#sec5sec4) across 128 threads
- * where each thread owns 4 consecutive keys. The final partitioning is striped.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockRadixSort for 128 threads owning 4 integer keys each
- * \brief Performs a radix sort across a [<em>blocked arrangement</em>](index.html#sec5sec4) of keys and values, leaving them in a [<em>striped arrangement</em>](index.html#sec5sec4).
- *
- * BlockRadixSort can only accommodate one associated tile of values. To "truck along"
- * more than one tile of values, simply perform a key-value sort of the keys paired
- * with a temporary value array that enumerates the key indices. The reordered indices
- * can then be used as a gather-vector for exchanging other associated tile data through
- * shared memory.
- *
- * \smemreuse
- *
- * The code snippet below illustrates a sort of 512 integer keys and values that
- * are initially partitioned in a [<em>blocked arrangement</em>](index.html#sec5sec4) across 128 threads
- * where each thread owns 4 consecutive pairs. The final partitioning is striped.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockRadixSort for 128 threads owning 4 integer keys and values each
- /// Never use a raking thread that will have no valid data (e.g., when BLOCK_THREADS is 62 and SEGMENT_LENGTH is 2, we should only use 31 raking threads)
- * The cub::BlockReduce class provides [<em>collective</em>](index.html#sec0) methods for computing a parallel reduction of items partitioned across a CUDA thread block.
- * \brief The BlockReduce class provides [<em>collective</em>](index.html#sec0) methods for computing a parallel reduction of items partitioned across a CUDA thread block. 
- * \ingroup BlockModule
- *
- * \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 list of input elements.
- *
- * \par
- * Optionally, BlockReduce can be specialized by algorithm to accommodate different latency/throughput workload profiles:
- * \tparam BLOCK_THREADS The thread block size in threads
- * \tparam ALGORITHM <b>[optional]</b> cub::BlockReduceAlgorithm enumerator specifying the underlying algorithm to use (default: cub::BLOCK_REDUCE_RAKING)
- *
- * \par Performance Considerations
- * - Very efficient (only one synchronization barrier).
- * - Zero bank conflicts for most types.
- * - Computation is slightly more efficient (i.e., having lower instruction overhead) for:
- * \brief Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockReduce()
- :
- temp_storage(PrivateStorage()),
- linear_tid(threadIdx.x)
- {}
-
-
- /**
- * \brief Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockReduce(
- TempStorage &temp_storage) ///< [in] Reference to memory allocation having layout type TempStorage
- :
- temp_storage(temp_storage.Alias()),
- linear_tid(threadIdx.x)
- {}
-
-
- /**
- * \brief Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier
- */
- __device__ __forceinline__ BlockReduce(
- int linear_tid) ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- :
- temp_storage(PrivateStorage()),
- linear_tid(linear_tid)
- {}
-
-
- /**
- * \brief Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier.
- */
- __device__ __forceinline__ BlockReduce(
- TempStorage &temp_storage, ///< [in] Reference to memory allocation having layout type TempStorage
- int linear_tid) ///< [in] <b>[optional]</b> A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- * \brief Computes a block-wide reduction for thread<sub>0</sub> using the specified binary reduction functor. Each thread contributes one input element.
- *
- * The return value is undefined in threads other than thread<sub>0</sub>.
- *
- * Supports non-commutative reduction operators.
- *
- * \smemreuse
- *
- * The code snippet below illustrates a max reduction of 128 integer items that
- * are partitioned across 128 threads.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockReduce for 128 threads on type int
- * \brief Computes a block-wide reduction for thread<sub>0</sub> using the specified binary reduction functor. Each thread contributes an array of consecutive input elements.
- *
- * The return value is undefined in threads other than thread<sub>0</sub>.
- *
- * Supports non-commutative reduction operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates a max reduction of 512 integer items that
- * are partitioned in a [<em>blocked arrangement</em>](index.html#sec5sec4) across 128 threads
- * where each thread owns 4 consecutive items.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockReduce for 128 threads on type int
- * \brief Computes a block-wide reduction for thread<sub>0</sub> using the specified binary reduction functor. The first \p num_valid threads each contribute one input element.
- *
- * The return value is undefined in threads other than thread<sub>0</sub>.
- *
- * Supports non-commutative reduction operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates a max reduction of a partially-full tile of integer items that
- * \brief Computes a block-wide reduction for thread<sub>0</sub> using addition (+) as the reduction operator. Each thread contributes one input element.
- *
- * The return value is undefined in threads other than thread<sub>0</sub>.
- *
- * \smemreuse
- *
- * The code snippet below illustrates a sum reduction of 128 integer items that
- * are partitioned across 128 threads.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockReduce for 128 threads on type int
- * \brief Computes a block-wide reduction for thread<sub>0</sub> using addition (+) as the reduction operator. Each thread contributes an array of consecutive input elements.
- *
- * The return value is undefined in threads other than thread<sub>0</sub>.
- *
- * \smemreuse
- *
- * The code snippet below illustrates a sum reduction of 512 integer items that
- * are partitioned in a [<em>blocked arrangement</em>](index.html#sec5sec4) across 128 threads
- * where each thread owns 4 consecutive items.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockReduce for 128 threads on type int
- * // Obtain a segment of consecutive items that are blocked across threads
- * int thread_data[4];
- * ...
- *
- * // Compute the block-wide sum for thread0
- * int aggregate = BlockReduce(temp_storage).Sum(thread_data);
- *
- * \endcode
- *
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- */
- template <int ITEMS_PER_THREAD>
- __device__ __forceinline__ T Sum(
- T (&inputs)[ITEMS_PER_THREAD]) ///< [in] Calling thread's input segment
- {
- // Reduce partials
- T partial = ThreadReduce(inputs, cub::Sum());
- return Sum(partial);
- }
-
-
- /**
- * \brief Computes a block-wide reduction for thread<sub>0</sub> using addition (+) as the reduction operator. The first \p num_valid threads each contribute one input element.
- *
- * The return value is undefined in threads other than thread<sub>0</sub>.
- *
- * \smemreuse
- *
- * The code snippet below illustrates a sum reduction of a partially-full tile of integer items that
- * The cub::BlockScan class provides [<em>collective</em>](index.html#sec0) methods for computing a parallel prefix sum/scan of items partitioned across a CUDA thread block.
- * \brief BlockScanAlgorithm enumerates alternative algorithms for cub::BlockScan to compute a parallel prefix scan across a CUDA thread block.
- */
-enum BlockScanAlgorithm
-{
-
- /**
- * \par Overview
- * An efficient "raking reduce-then-scan" prefix scan algorithm. Execution is comprised of five phases:
- * -# Upsweep sequential reduction in registers (if threads contribute more than one input each). Each thread then places the partial reduction of its item(s) into shared memory.
- * -# Upsweep sequential reduction in shared memory. Threads within a single warp rake across segments of shared partial reductions.
- * -# A warp-synchronous Kogge-Stone style exclusive scan within the raking warp.
- * -# Downsweep sequential exclusive scan in shared memory. Threads within a single warp rake across segments of shared partial reductions, seeded with the warp-scan output.
- * -# Downsweep sequential scan in registers (if threads contribute more than one input), seeded with the raking scan output.
- *
- * \par
- * \image html block_scan_raking.png
- * <div class="centercaption">\p BLOCK_SCAN_RAKING data flow for a hypothetical 16-thread threadblock and 4-thread raking warp.</div>
- *
- * \par Performance Considerations
- * - Although this variant may suffer longer turnaround latencies when the
- * GPU is under-occupied, it can often provide higher overall throughput
- * across the GPU when suitably occupied.
- */
- BLOCK_SCAN_RAKING,
-
-
- /**
- * \par Overview
- * Similar to cub::BLOCK_SCAN_RAKING, but with fewer shared memory reads at
- * the expense of higher register pressure. Raking threads preserve their
- * "upsweep" segment of values in registers while performing warp-synchronous
- * scan, allowing the "downsweep" not to re-read them from shared memory.
- */
- BLOCK_SCAN_RAKING_MEMOIZE,
-
-
- /**
- * \par Overview
- * A quick "tiled warpscans" prefix scan algorithm. Execution is comprised of four phases:
- * -# Upsweep sequential reduction in registers (if threads contribute more than one input each). Each thread then places the partial reduction of its item(s) into shared memory.
- * -# Compute a shallow, but inefficient warp-synchronous Kogge-Stone style scan within each warp.
- * -# A propagation phase where the warp scan outputs in each warp are updated with the aggregate from each preceding warp.
- * -# Downsweep sequential scan in registers (if threads contribute more than one input), seeded with the raking scan output.
- *
- * \par
- * \image html block_scan_warpscans.png
- * <div class="centercaption">\p BLOCK_SCAN_WARP_SCANS data flow for a hypothetical 16-thread threadblock and 4-thread raking warp.</div>
- *
- * \par Performance Considerations
- * - Although this variant may suffer lower overall throughput across the
- * GPU because due to a heavy reliance on inefficient warpscans, it can
- * often provide lower turnaround latencies when the GPU is under-occupied.
- * \brief The BlockScan class provides [<em>collective</em>](index.html#sec0) methods for computing a parallel prefix sum/scan of items partitioned across a CUDA thread block. 
- * \ingroup BlockModule
- *
- * \par Overview
- * Given a list of input elements and a binary reduction operator, a [<em>prefix scan</em>](http://en.wikipedia.org/wiki/Prefix_sum)
- * produces an output list where each element is computed to be the reduction
- * of the elements occurring earlier in the input list. <em>Prefix sum</em>
- * connotes a prefix scan with the addition operator. The term \em inclusive indicates
- * that the <em>i</em><sup>th</sup> output reduction incorporates the <em>i</em><sup>th</sup> input.
- * The term \em exclusive indicates the <em>i</em><sup>th</sup> input is not incorporated into
- * the <em>i</em><sup>th</sup> output reduction.
- *
- * \par
- * Optionally, BlockScan can be specialized by algorithm to accommodate different latency/throughput workload profiles:
- * \tparam BLOCK_THREADS The thread block size in threads
- * \tparam ALGORITHM <b>[optional]</b> cub::BlockScanAlgorithm enumerator specifying the underlying algorithm to use (default: cub::BLOCK_SCAN_RAKING)
- *
- * \par A Simple Example
- * \blockcollective{BlockScan}
- * \par
- * The code snippet below illustrates an exclusive prefix sum of 512 integer items that
- * are partitioned in a [<em>blocked arrangement</em>](index.html#sec5sec4) across 128 threads
- * where each thread owns 4 consecutive items.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockScan for 128 threads on type int
- * \brief Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockScan()
- :
- temp_storage(PrivateStorage()),
- linear_tid(threadIdx.x)
- {}
-
-
- /**
- * \brief Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockScan(
- TempStorage &temp_storage) ///< [in] Reference to memory allocation having layout type TempStorage
- :
- temp_storage(temp_storage.Alias()),
- linear_tid(threadIdx.x)
- {}
-
-
- /**
- * \brief Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier
- */
- __device__ __forceinline__ BlockScan(
- int linear_tid) ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- :
- temp_storage(PrivateStorage()),
- linear_tid(linear_tid)
- {}
-
-
- /**
- * \brief Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier.
- */
- __device__ __forceinline__ BlockScan(
- TempStorage &temp_storage, ///< [in] Reference to memory allocation having layout type TempStorage
- int linear_tid) ///< [in] <b>[optional]</b> A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- * \brief Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates an exclusive prefix sum of 128 integer items that
- * are partitioned across 128 threads.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockScan for 128 threads on type int
- * \brief Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Instead of using 0 as the block-wide prefix, the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * The \p block_prefix_op functor must implement a member function <tt>T operator()(T block_aggregate)</tt>.
- * The functor's input parameter \p block_aggregate is the same value also returned by the scan operation.
- * The functor will be invoked by the first warp of threads in the block, however only the return value from
- * <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates a single thread block that progressively
- * computes an exclusive prefix sum over multiple "tiles" of input using a
- * prefix functor to maintain a running total between block-wide scans. Each tile consists
- * of 128 integer items that are partitioned across 128 threads.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * // A stateful callback functor that maintains a running prefix to be applied
- * Suppose the input \p d_data is <tt>1, 1, 1, 1, 1, 1, 1, 1, ...</tt>.
- * The corresponding output for the first segment will be <tt>0, 1, ..., 127</tt>.
- * The output for the second segment will be <tt>128, 129, ..., 255</tt>. Furthermore,
- * the value \p 128 will be stored in \p block_aggregate for all threads after each scan.
- *
- * \tparam BlockPrefixOp <b>[inferred]</b> Call-back functor type having member <tt>T operator()(T block_aggregate)</tt>
- */
- template <typename BlockPrefixOp>
- __device__ __forceinline__ void ExclusiveSum(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- T &block_aggregate, ///< [out] block-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs.
- * \brief Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates an exclusive prefix sum of 512 integer items that
- * are partitioned in a [<em>blocked arrangement</em>](index.html#sec5sec4) across 128 threads
- * where each thread owns 4 consecutive items.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockScan for 128 threads on type int
- * \brief Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates an exclusive prefix sum of 512 integer items that
- * are partitioned in a [<em>blocked arrangement</em>](index.html#sec5sec4) across 128 threads
- * where each thread owns 4 consecutive items.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockScan for 128 threads on type int
- * \brief Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Instead of using 0 as the block-wide prefix, the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * The \p block_prefix_op functor must implement a member function <tt>T operator()(T block_aggregate)</tt>.
- * The functor's input parameter \p block_aggregate is the same value also returned by the scan operation.
- * The functor will be invoked by the first warp of threads in the block, however only the return value from
- * <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates a single thread block that progressively
- * computes an exclusive prefix sum over multiple "tiles" of input using a
- * prefix functor to maintain a running total between block-wide scans. Each tile consists
- * of 512 integer items that are partitioned in a [<em>blocked arrangement</em>](index.html#sec5sec4)
- * across 128 threads where each thread owns 4 consecutive items.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * // A stateful callback functor that maintains a running prefix to be applied
- * Suppose the input \p d_data is <tt>1, 1, 1, 1, 1, 1, 1, 1, ...</tt>.
- * The corresponding output for the first segment will be <tt>0, 1, 2, 3, ..., 510, 511</tt>.
- * The output for the second segment will be <tt>512, 513, 514, 515, ..., 1022, 1023</tt>. Furthermore,
- * the value \p 512 will be stored in \p block_aggregate for all threads after each scan.
- *
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam BlockPrefixOp <b>[inferred]</b> Call-back functor type having member <tt>T operator()(T block_aggregate)</tt>
- */
- template <
- int ITEMS_PER_THREAD,
- typename BlockPrefixOp>
- __device__ __forceinline__ void ExclusiveSum(
- T (&input)[ITEMS_PER_THREAD], ///< [in] Calling thread's input items
- T (&output)[ITEMS_PER_THREAD], ///< [out] Calling thread's output items (may be aliased to \p input)
- T &block_aggregate, ///< [out] block-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs.
- {
- // Reduce consecutive thread items in registers
- Sum scan_op;
- T thread_partial = ThreadReduce(input, scan_op);
- * \brief Computes an exclusive block-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * Supports non-commutative scan operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates an exclusive prefix max scan of 128 integer items that
- * are partitioned across 128 threads.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockScan for 128 threads on type int
- * \brief Computes an exclusive block-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * The \p block_prefix_op functor must implement a member function <tt>T operator()(T block_aggregate)</tt>.
- * The functor's input parameter \p block_aggregate is the same value also returned by the scan operation.
- * The functor will be invoked by the first warp of threads in the block, however only the return value from
- * <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.
- *
- * Supports non-commutative scan operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates a single thread block that progressively
- * computes an exclusive prefix max scan over multiple "tiles" of input using a
- * prefix functor to maintain a running total between block-wide scans. Each tile consists
- * of 128 integer items that are partitioned across 128 threads.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * // A stateful callback functor that maintains a running prefix to be applied
- * Suppose the input \p d_data is <tt>0, -1, 2, -3, 4, -5, ...</tt>.
- * The corresponding output for the first segment will be <tt>INT_MIN, 0, 0, 2, ..., 124, 126</tt>.
- * The output for the second segment will be <tt>126, 128, 128, 130, ..., 252, 254</tt>. Furthermore,
- * \p block_aggregate will be assigned \p 126 in all threads after the first scan, assigned \p 254 after the second
- * scan, etc.
- *
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- * \tparam BlockPrefixOp <b>[inferred]</b> Call-back functor type having member <tt>T operator()(T block_aggregate)</tt>
- */
- template <
- typename ScanOp,
- typename BlockPrefixOp>
- __device__ __forceinline__ void ExclusiveScan(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- T identity, ///< [in] Identity value
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate, ///< [out] block-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs.
- * \brief Computes an exclusive block-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes an array of consecutive input elements.
- *
- * Supports non-commutative scan operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates an exclusive prefix max scan of 512 integer items that
- * are partitioned in a [<em>blocked arrangement</em>](index.html#sec5sec4) across 128 threads
- * where each thread owns 4 consecutive items.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockScan for 128 threads on type int
- * \brief Computes an exclusive block-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * Supports non-commutative scan operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates an exclusive prefix max scan of 512 integer items that
- * are partitioned in a [<em>blocked arrangement</em>](index.html#sec5sec4) across 128 threads
- * where each thread owns 4 consecutive items.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockScan for 128 threads on type int
- * \brief Computes an exclusive block-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes an array of consecutive input elements. the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * The \p block_prefix_op functor must implement a member function <tt>T operator()(T block_aggregate)</tt>.
- * The functor's input parameter \p block_aggregate is the same value also returned by the scan operation.
- * The functor will be invoked by the first warp of threads in the block, however only the return value from
- * <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.
- *
- * Supports non-commutative scan operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates a single thread block that progressively
- * computes an exclusive prefix max scan over multiple "tiles" of input using a
- * prefix functor to maintain a running total between block-wide scans. Each tile consists
- * of 128 integer items that are partitioned across 128 threads.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * // A stateful callback functor that maintains a running prefix to be applied
- * Suppose the input \p d_data is <tt>0, -1, 2, -3, 4, -5, ...</tt>.
- * The corresponding output for the first segment will be <tt>INT_MIN, 0, 0, 2, 2, 4, ..., 508, 510</tt>.
- * The output for the second segment will be <tt>510, 512, 512, 514, 514, 516, ..., 1020, 1022</tt>. Furthermore,
- * \p block_aggregate will be assigned \p 510 in all threads after the first scan, assigned \p 1022 after the second
- * scan, etc.
- *
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- * \tparam BlockPrefixOp <b>[inferred]</b> Call-back functor type having member <tt>T operator()(T block_aggregate)</tt>
- */
- template <
- int ITEMS_PER_THREAD,
- typename ScanOp,
- typename BlockPrefixOp>
- __device__ __forceinline__ void ExclusiveScan(
- T (&input)[ITEMS_PER_THREAD], ///< [in] Calling thread's input items
- T (&output)[ITEMS_PER_THREAD], ///< [out] Calling thread's output items (may be aliased to \p input)
- T identity, ///< [in] Identity value
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate, ///< [out] block-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs.
- {
- // Reduce consecutive thread items in registers
- T thread_partial = ThreadReduce(input, scan_op);
- * \brief Computes an exclusive block-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. With no identity value, the output computed for <em>thread</em><sub>0</sub> is undefined.
- *
- * Supports non-commutative scan operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- */
- template <typename ScanOp>
- __device__ __forceinline__ void ExclusiveScan(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- * \brief Computes an exclusive block-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. Also provides every thread with the block-wide \p block_aggregate of all inputs. With no identity value, the output computed for <em>thread</em><sub>0</sub> is undefined.
- *
- * Supports non-commutative scan operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- */
- template <typename ScanOp>
- __device__ __forceinline__ void ExclusiveScan(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate) ///< [out] block-wide aggregate reduction of input items
- * \brief Computes an exclusive block-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * The \p block_prefix_op functor must implement a member function <tt>T operator()(T block_aggregate)</tt>.
- * The functor's input parameter \p block_aggregate is the same value also returned by the scan operation.
- * The functor will be invoked by the first warp of threads in the block, however only the return value from
- * <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.
- *
- * Supports non-commutative scan operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- * \tparam BlockPrefixOp <b>[inferred]</b> Call-back functor type having member <tt>T operator()(T block_aggregate)</tt>
- */
- template <
- typename ScanOp,
- typename BlockPrefixOp>
- __device__ __forceinline__ void ExclusiveScan(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate, ///< [out] block-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs.
- * \brief Computes an exclusive block-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes an array of consecutive input elements. With no identity value, the output computed for <em>thread</em><sub>0</sub> is undefined.
- *
- * Supports non-commutative scan operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- */
- template <
- int ITEMS_PER_THREAD,
- typename ScanOp>
- __device__ __forceinline__ void ExclusiveScan(
- T (&input)[ITEMS_PER_THREAD], ///< [in] Calling thread's input items
- T (&output)[ITEMS_PER_THREAD], ///< [out] Calling thread's output items (may be aliased to \p input)
- ScanOp scan_op) ///< [in] Binary scan operator
- {
- // Reduce consecutive thread items in registers
- T thread_partial = ThreadReduce(input, scan_op);
- * \brief Computes an exclusive block-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide \p block_aggregate of all inputs. With no identity value, the output computed for <em>thread</em><sub>0</sub> is undefined.
- *
- * Supports non-commutative scan operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- */
- template <
- int ITEMS_PER_THREAD,
- typename ScanOp>
- __device__ __forceinline__ void ExclusiveScan(
- T (&input)[ITEMS_PER_THREAD], ///< [in] Calling thread's input items
- T (&output)[ITEMS_PER_THREAD], ///< [out] Calling thread's output items (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate) ///< [out] block-wide aggregate reduction of input items
- {
- // Reduce consecutive thread items in registers
- T thread_partial = ThreadReduce(input, scan_op);
- * \brief Computes an exclusive block-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes an array of consecutive input elements. the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * The \p block_prefix_op functor must implement a member function <tt>T operator()(T block_aggregate)</tt>.
- * The functor's input parameter \p block_aggregate is the same value also returned by the scan operation.
- * The functor will be invoked by the first warp of threads in the block, however only the return value from
- * <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.
- *
- * Supports non-commutative scan operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- * \tparam BlockPrefixOp <b>[inferred]</b> Call-back functor type having member <tt>T operator()(T block_aggregate)</tt>
- */
- template <
- int ITEMS_PER_THREAD,
- typename ScanOp,
- typename BlockPrefixOp>
- __device__ __forceinline__ void ExclusiveScan(
- T (&input)[ITEMS_PER_THREAD], ///< [in] Calling thread's input items
- T (&output)[ITEMS_PER_THREAD], ///< [out] Calling thread's output items (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate, ///< [out] block-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs.
- {
- // Reduce consecutive thread items in registers
- T thread_partial = ThreadReduce(input, scan_op);
- * \brief Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates an inclusive prefix sum of 128 integer items that
- * are partitioned across 128 threads.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockScan for 128 threads on type int
- * \brief Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Instead of using 0 as the block-wide prefix, the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * The \p block_prefix_op functor must implement a member function <tt>T operator()(T block_aggregate)</tt>.
- * The functor's input parameter \p block_aggregate is the same value also returned by the scan operation.
- * The functor will be invoked by the first warp of threads in the block, however only the return value from
- * <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates a single thread block that progressively
- * computes an inclusive prefix sum over multiple "tiles" of input using a
- * prefix functor to maintain a running total between block-wide scans. Each tile consists
- * of 128 integer items that are partitioned across 128 threads.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * // A stateful callback functor that maintains a running prefix to be applied
- * Suppose the input \p d_data is <tt>1, 1, 1, 1, 1, 1, 1, 1, ...</tt>.
- * The corresponding output for the first segment will be <tt>1, 2, ..., 128</tt>.
- * The output for the second segment will be <tt>129, 130, ..., 256</tt>. Furthermore,
- * the value \p 128 will be stored in \p block_aggregate for all threads after each scan.
- *
- * \tparam BlockPrefixOp <b>[inferred]</b> Call-back functor type having member <tt>T operator()(T block_aggregate)</tt>
- */
- template <typename BlockPrefixOp>
- __device__ __forceinline__ void InclusiveSum(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- T &block_aggregate, ///< [out] block-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs.
- * \brief Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates an inclusive prefix sum of 512 integer items that
- * are partitioned in a [<em>blocked arrangement</em>](index.html#sec5sec4) across 128 threads
- * where each thread owns 4 consecutive items.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockScan for 128 threads on type int
- * \brief Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates an inclusive prefix sum of 512 integer items that
- * are partitioned in a [<em>blocked arrangement</em>](index.html#sec5sec4) across 128 threads
- * where each thread owns 4 consecutive items.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockScan for 128 threads on type int
- * \brief Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Instead of using 0 as the block-wide prefix, the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * The \p block_prefix_op functor must implement a member function <tt>T operator()(T block_aggregate)</tt>.
- * The functor's input parameter \p block_aggregate is the same value also returned by the scan operation.
- * The functor will be invoked by the first warp of threads in the block, however only the return value from
- * <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates a single thread block that progressively
- * computes an inclusive prefix sum over multiple "tiles" of input using a
- * prefix functor to maintain a running total between block-wide scans. Each tile consists
- * of 512 integer items that are partitioned in a [<em>blocked arrangement</em>](index.html#sec5sec4)
- * across 128 threads where each thread owns 4 consecutive items.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * // A stateful callback functor that maintains a running prefix to be applied
- * Suppose the input \p d_data is <tt>1, 1, 1, 1, 1, 1, 1, 1, ...</tt>.
- * The corresponding output for the first segment will be <tt>1, 2, 3, 4, ..., 511, 512</tt>.
- * The output for the second segment will be <tt>513, 514, 515, 516, ..., 1023, 1024</tt>. Furthermore,
- * the value \p 512 will be stored in \p block_aggregate for all threads after each scan.
- *
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam BlockPrefixOp <b>[inferred]</b> Call-back functor type having member <tt>T operator()(T block_aggregate)</tt>
- */
- template <
- int ITEMS_PER_THREAD,
- typename BlockPrefixOp>
- __device__ __forceinline__ void InclusiveSum(
- T (&input)[ITEMS_PER_THREAD], ///< [in] Calling thread's input items
- T (&output)[ITEMS_PER_THREAD], ///< [out] Calling thread's output items (may be aliased to \p input)
- T &block_aggregate, ///< [out] block-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs.
- * \brief Computes an inclusive block-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * Supports non-commutative scan operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates an inclusive prefix max scan of 128 integer items that
- * are partitioned across 128 threads.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockScan for 128 threads on type int
- * \brief Computes an inclusive block-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * The \p block_prefix_op functor must implement a member function <tt>T operator()(T block_aggregate)</tt>.
- * The functor's input parameter \p block_aggregate is the same value also returned by the scan operation.
- * The functor will be invoked by the first warp of threads in the block, however only the return value from
- * <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.
- *
- * Supports non-commutative scan operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates a single thread block that progressively
- * computes an inclusive prefix max scan over multiple "tiles" of input using a
- * prefix functor to maintain a running total between block-wide scans. Each tile consists
- * of 128 integer items that are partitioned across 128 threads.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * // A stateful callback functor that maintains a running prefix to be applied
- * Suppose the input \p d_data is <tt>0, -1, 2, -3, 4, -5, ...</tt>.
- * The corresponding output for the first segment will be <tt>0, 0, 2, 2, ..., 126, 126</tt>.
- * The output for the second segment will be <tt>128, 128, 130, 130, ..., 254, 254</tt>. Furthermore,
- * \p block_aggregate will be assigned \p 126 in all threads after the first scan, assigned \p 254 after the second
- * scan, etc.
- *
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- * \tparam BlockPrefixOp <b>[inferred]</b> Call-back functor type having member <tt>T operator()(T block_aggregate)</tt>
- */
- template <
- typename ScanOp,
- typename BlockPrefixOp>
- __device__ __forceinline__ void InclusiveScan(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate, ///< [out] block-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs.
- * \brief Computes an inclusive block-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes an array of consecutive input elements.
- *
- * Supports non-commutative scan operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates an inclusive prefix max scan of 512 integer items that
- * are partitioned in a [<em>blocked arrangement</em>](index.html#sec5sec4) across 128 threads
- * where each thread owns 4 consecutive items.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockScan for 128 threads on type int
- * \brief Computes an inclusive block-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * Supports non-commutative scan operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates an inclusive prefix max scan of 512 integer items that
- * are partitioned in a [<em>blocked arrangement</em>](index.html#sec5sec4) across 128 threads
- * where each thread owns 4 consecutive items.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize BlockScan for 128 threads on type int
- * \brief Computes an inclusive block-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes an array of consecutive input elements. the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- *
- * The \p block_prefix_op functor must implement a member function <tt>T operator()(T block_aggregate)</tt>.
- * The functor's input parameter \p block_aggregate is the same value also returned by the scan operation.
- * The functor will be invoked by the first warp of threads in the block, however only the return value from
- * <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.
- *
- * Supports non-commutative scan operators.
- *
- * \blocked
- *
- * \smemreuse
- *
- * The code snippet below illustrates a single thread block that progressively
- * computes an inclusive prefix max scan over multiple "tiles" of input using a
- * prefix functor to maintain a running total between block-wide scans. Each tile consists
- * of 128 integer items that are partitioned across 128 threads.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * // A stateful callback functor that maintains a running prefix to be applied
- * Suppose the input \p d_data is <tt>0, -1, 2, -3, 4, -5, ...</tt>.
- * The corresponding output for the first segment will be <tt>0, 0, 2, 2, 4, 4, ..., 510, 510</tt>.
- * The output for the second segment will be <tt>512, 512, 514, 514, 516, 516, ..., 1022, 1022</tt>. Furthermore,
- * \p block_aggregate will be assigned \p 510 in all threads after the first scan, assigned \p 1022 after the second
- * scan, etc.
- *
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- * \tparam BlockPrefixOp <b>[inferred]</b> Call-back functor type having member <tt>T operator()(T block_aggregate)</tt>
- */
- template <
- int ITEMS_PER_THREAD,
- typename ScanOp,
- typename BlockPrefixOp>
- __device__ __forceinline__ void InclusiveScan(
- T (&input)[ITEMS_PER_THREAD], ///< [in] Calling thread's input items
- T (&output)[ITEMS_PER_THREAD], ///< [out] Calling thread's output items (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate, ///< [out] block-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs.
- * \tparam T <b>[inferred]</b> The data type to store.
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam OutputIteratorRA <b>[inferred]</b> The random-access iterator type for output (may be a simple pointer type).
- */
-template <
- PtxStoreModifier MODIFIER,
- typename T,
- int ITEMS_PER_THREAD,
- typename OutputIteratorRA>
-__device__ __forceinline__ void StoreBlocked(
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- OutputIteratorRA block_itr, ///< [in] The thread block's base output iterator for storing to
- T (&items)[ITEMS_PER_THREAD]) ///< [in] Data to store
-{
- // Store directly in thread-blocked order
- #pragma unroll
- for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
- * \brief Store a blocked arrangement of items across a thread block into a linear segment of items using the specified cache modifier, guarded by range
- * \tparam T <b>[inferred]</b> The data type to store.
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam OutputIteratorRA <b>[inferred]</b> The random-access iterator type for output (may be a simple pointer type).
- */
-template <
- PtxStoreModifier MODIFIER,
- typename T,
- int ITEMS_PER_THREAD,
- typename OutputIteratorRA>
-__device__ __forceinline__ void StoreBlocked(
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- OutputIteratorRA block_itr, ///< [in] The thread block's base output iterator for storing to
- T (&items)[ITEMS_PER_THREAD], ///< [in] Data to store
- int valid_items) ///< [in] Number of valid items to write
-{
- // Store directly in thread-blocked order
- #pragma unroll
- for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
- {
- if (ITEM + (linear_tid * ITEMS_PER_THREAD) < valid_items)
- * \tparam BLOCK_THREADS The thread block size in threads
- * \tparam T <b>[inferred]</b> The data type to store.
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam OutputIteratorRA <b>[inferred]</b> The random-access iterator type for output (may be a simple pointer type).
- */
-template <
- PtxStoreModifier MODIFIER,
- int BLOCK_THREADS,
- typename T,
- int ITEMS_PER_THREAD,
- typename OutputIteratorRA>
-__device__ __forceinline__ void StoreStriped(
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- OutputIteratorRA block_itr, ///< [in] The thread block's base output iterator for storing to
- T (&items)[ITEMS_PER_THREAD]) ///< [in] Data to store
-{
- // Store directly in striped order
- #pragma unroll
- for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
- * \brief Store a striped arrangement of data across the thread block into a linear segment of items using the specified cache modifier, guarded by range
- * \tparam BLOCK_THREADS The thread block size in threads
- * \tparam T <b>[inferred]</b> The data type to store.
- * \tparam ITEMS_PER_THREAD <b>[inferred]</b> The number of consecutive items partitioned onto each thread.
- * \tparam OutputIteratorRA <b>[inferred]</b> The random-access iterator type for output (may be a simple pointer type).
- */
-template <
- PtxStoreModifier MODIFIER,
- int BLOCK_THREADS,
- typename T,
- int ITEMS_PER_THREAD,
- typename OutputIteratorRA>
-__device__ __forceinline__ void StoreStriped(
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- OutputIteratorRA block_itr, ///< [in] The thread block's base output iterator for storing to
- T (&items)[ITEMS_PER_THREAD], ///< [in] Data to store
- int valid_items) ///< [in] Number of valid items to write
-{
- // Store directly in striped order
- #pragma unroll
- for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
- {
- if ((ITEM * BLOCK_THREADS) + linear_tid < valid_items)
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- OutputIteratorRA block_itr, ///< [in] The thread block's base output iterator for storing to
- T (&items)[ITEMS_PER_THREAD]) ///< [out] Data to load
-{
- int tid = linear_tid & (PtxArchProps::WARP_THREADS - 1);
- int wid = linear_tid >> PtxArchProps::LOG_WARP_THREADS;
- int warp_offset = wid * PtxArchProps::WARP_THREADS * ITEMS_PER_THREAD;
-
- // Store directly in warp-striped order
- #pragma unroll
- for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
- * \brief Store a warp-striped arrangement of data across the thread block into a linear segment of items using the specified cache modifier, guarded by range
- *
- * \warpstriped
- *
- * \par Usage Considerations
- * The number of threads in the thread block must be a multiple of the architecture's warp size.
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- OutputIteratorRA block_itr, ///< [in] The thread block's base output iterator for storing to
- T (&items)[ITEMS_PER_THREAD], ///< [in] Data to store
- int valid_items) ///< [in] Number of valid items to write
-{
- int tid = linear_tid & (PtxArchProps::WARP_THREADS - 1);
- int wid = linear_tid >> PtxArchProps::LOG_WARP_THREADS;
- int warp_offset = wid * PtxArchProps::WARP_THREADS * ITEMS_PER_THREAD;
-
- // Store directly in warp-striped order
- #pragma unroll
- for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
- {
- if (warp_offset + tid + (ITEM * PtxArchProps::WARP_THREADS) < valid_items)
- int linear_tid, ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- T *block_ptr, ///< [in] Input pointer for storing from
- T (&items)[ITEMS_PER_THREAD]) ///< [in] Data to store
-{
- enum
- {
- // Maximum CUDA vector size is 4 elements
- MAX_VEC_SIZE = CUB_MIN(4, ITEMS_PER_THREAD),
-
- // Vector size must be a power of two and an even divisor of the items per thread
- * \brief cub::BlockStoreAlgorithm enumerates alternative algorithms for cub::BlockStore to write a blocked arrangement of items across a CUDA thread block to a linear segment of memory.
- */
-enum BlockStoreAlgorithm
-{
- /**
- * \par Overview
- *
- * A [<em>blocked arrangement</em>](index.html#sec5sec4) of data is written
- * directly to memory. The thread block writes items in a parallel "raking" fashion:
- * thread<sub><em>i</em></sub> writes the <em>i</em><sup>th</sup> segment of consecutive elements.
- *
- * \par Performance Considerations
- * - The utilization of memory transactions (coalescing) decreases as the
- * access stride between threads increases (i.e., the number items per thread).
- */
- BLOCK_STORE_DIRECT,
-
- /**
- * \par Overview
- *
- * A [<em>blocked arrangement</em>](index.html#sec5sec4) of data is written directly
- * to memory using CUDA's built-in vectorized stores as a coalescing optimization.
- * The thread block writes items in a parallel "raking" fashion: thread<sub><em>i</em></sub> uses vector stores to
- * write the <em>i</em><sup>th</sup> segment of consecutive elements.
- *
- * For example, <tt>st.global.v4.s32</tt> instructions will be generated when \p T = \p int and \p ITEMS_PER_THREAD > 4.
- *
- * \par Performance Considerations
- * - The utilization of memory transactions (coalescing) remains high until the the
- * access stride between threads (i.e., the number items per thread) exceeds the
- * maximum vector store width (typically 4 items or 64B, whichever is lower).
- * - The following conditions will prevent vectorization and writing will fall back to cub::BLOCK_STORE_DIRECT:
- * - \p ITEMS_PER_THREAD is odd
- * - The \p OutputIteratorRA is not a simple pointer type
- * - The block output offset is not quadword-aligned
- * - The data type \p T is not a built-in primitive or CUDA vector type (e.g., \p short, \p int2, \p double, \p float2, etc.)
- */
- BLOCK_STORE_VECTORIZE,
-
- /**
- * \par Overview
- * A [<em>blocked arrangement</em>](index.html#sec5sec4) is locally
- * transposed into a [<em>striped arrangement</em>](index.html#sec5sec4)
- * which is then written to memory. More specifically, cub::BlockExchange
- * used to locally reorder the items into a
- * [<em>striped arrangement</em>](index.html#sec5sec4), after which the
- * thread block writes items in a parallel "strip-mining" fashion: consecutive
- * items owned by thread<sub><em>i</em></sub> are written to memory with
- * stride \p BLOCK_THREADS between them.
- *
- * \par Performance Considerations
- * - The utilization of memory transactions (coalescing) remains high regardless
- * of items written per thread.
- * - The local reordering incurs slightly longer latencies and throughput than the
- * direct cub::BLOCK_STORE_DIRECT and cub::BLOCK_STORE_VECTORIZE alternatives.
- */
- BLOCK_STORE_TRANSPOSE,
-
- /**
- * \par Overview
- * A [<em>blocked arrangement</em>](index.html#sec5sec4) is locally
- * transposed into a [<em>warp-striped arrangement</em>](index.html#sec5sec4)
- * which is then written to memory. More specifically, cub::BlockExchange used
- * to locally reorder the items into a
- * [<em>warp-striped arrangement</em>](index.html#sec5sec4), after which
- * each warp writes its own contiguous segment in a parallel "strip-mining" fashion:
- * consecutive items owned by lane<sub><em>i</em></sub> are written to memory
- * with stride \p WARP_THREADS between them.
- *
- * \par Performance Considerations
- * - The utilization of memory transactions (coalescing) remains high regardless
- * of items written per thread.
- * - The local reordering incurs slightly longer latencies and throughput than the
- * direct cub::BLOCK_STORE_DIRECT and cub::BLOCK_STORE_VECTORIZE alternatives.
- */
- BLOCK_STORE_WARP_TRANSPOSE,
-};
-
-
-
-/**
- * \addtogroup BlockModule
- * @{
- */
-
-
-/**
- * \brief The BlockStore class provides [<em>collective</em>](index.html#sec0) data movement methods for writing a [<em>blocked arrangement</em>](index.html#sec5sec4) of items partitioned across a CUDA thread block to a linear segment of memory. 
- *
- * \par Overview
- * The BlockStore class provides a single data movement abstraction that can be specialized
- * to implement different cub::BlockStoreAlgorithm strategies. This facilitates different
- * performance policies for different architectures, data types, granularity sizes, etc.
- *
- * \par Optionally, BlockStore can be specialized by different data movement strategies:
- * -# <b>cub::BLOCK_STORE_DIRECT</b>. A [<em>blocked arrangement</em>](index.html#sec5sec4) of data is written
- * directly to memory. [More...](\ref cub::BlockStoreAlgorithm)
- * -# <b>cub::BLOCK_STORE_VECTORIZE</b>. A [<em>blocked arrangement</em>](index.html#sec5sec4)
- * of data is written directly to memory using CUDA's built-in vectorized stores as a
- * \tparam WARP_TIME_SLICING <b>[optional]</b> For transposition-based cub::BlockStoreAlgorithm parameterizations that utilize shared memory: When \p true, only use enough shared memory for a single warp's worth of data, time-slicing the block-wide exchange over multiple synchronized rounds (default: false)
- *
- * \par A Simple Example
- * \blockcollective{BlockStore}
- * \par
- * The code snippet below illustrates the storing of a "blocked" arrangement
- * of 512 integers across 128 threads (where each thread owns 4 consecutive items)
- * into a linear segment of memory. The store is specialized for \p BLOCK_STORE_WARP_TRANSPOSE,
- * meaning items are locally reordered among threads so that memory references will be
- * efficiently coalesced using a warp-striped access pattern.
- * \brief Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockStore()
- :
- temp_storage(PrivateStorage()),
- linear_tid(threadIdx.x)
- {}
-
-
- /**
- * \brief Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ BlockStore(
- TempStorage &temp_storage) ///< [in] Reference to memory allocation having layout type TempStorage
- :
- temp_storage(temp_storage.Alias()),
- linear_tid(threadIdx.x)
- {}
-
-
- /**
- * \brief Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier
- */
- __device__ __forceinline__ BlockStore(
- int linear_tid) ///< [in] A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- :
- temp_storage(PrivateStorage()),
- linear_tid(linear_tid)
- {}
-
-
- /**
- * \brief Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier.
- */
- __device__ __forceinline__ BlockStore(
- TempStorage &temp_storage, ///< [in] Reference to memory allocation having layout type TempStorage
- int linear_tid) ///< [in] <b>[optional]</b> A suitable 1D thread-identifier for the calling thread (e.g., <tt>(threadIdx.y * blockDim.x) + linear_tid</tt> for 2D thread blocks)
- * The cub::BlockHistogramAtomic class provides atomic-based methods for constructing block-wide histograms from data samples partitioned across a CUDA thread block.
- */
-
-#pragma once
-
-#include "../../util_namespace.cuh"
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-
-/**
- * \brief The BlockHistogramAtomic class provides atomic-based methods for constructing block-wide histograms from data samples partitioned across a CUDA thread block.
- * The cub::BlockHistogramSort class provides sorting-based methods for constructing block-wide histograms from data samples partitioned across a CUDA thread block.
- */
-
-#pragma once
-
-#include "../../block/block_radix_sort.cuh"
-#include "../../block/block_discontinuity.cuh"
-#include "../../util_namespace.cuh"
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-
-
-/**
- * \brief The BlockHistogramSort class provides sorting-based methods for constructing block-wide histograms from data samples partitioned across a CUDA thread block.
- */
-template <
- typename T,
- int BLOCK_THREADS,
- int ITEMS_PER_THREAD,
- int BINS>
-struct BlockHistogramSort
-{
- // Parameterize BlockRadixSort type for our thread block
- /// 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)
- /// 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 input, ///< [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;
- unsigned int warp_offset = warp_id * LOGICAL_WARP_SIZE;
- unsigned int warp_num_valid = (FULL_TILE && EVEN_WARP_MULTIPLE) ?
- /// 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 input, ///< [in] Calling thread's input partial reductions
- int num_valid, ///< [in] Number of valid elements (may be less than BLOCK_THREADS)
- /// Computes an exclusive threadblock-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- template <typename ScanOp>
- __device__ __forceinline__ void ExclusiveScan(
- T input, ///< [in] Calling thread's input items
- T &output, ///< [out] Calling thread's output items (may be aliased to \p input)
- const T &identity, ///< [in] Identity value
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate) ///< [out] Threadblock-wide aggregate reduction of input items
- /// Computes an exclusive threadblock-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- template <
- typename ScanOp,
- typename BlockPrefixOp>
- __device__ __forceinline__ void ExclusiveScan(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- T identity, ///< [in] Identity value
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate, ///< [out] Threadblock-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a threadblock-wide prefix to be applied to all inputs.
- /// Computes an exclusive threadblock-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. Also provides every thread with the block-wide \p block_aggregate of all inputs. With no identity value, the output computed for <em>thread</em><sub>0</sub> is undefined.
- template <typename ScanOp>
- __device__ __forceinline__ void ExclusiveScan(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate) ///< [out] Threadblock-wide aggregate reduction of input items
- /// Computes an exclusive threadblock-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- template <
- typename ScanOp,
- typename BlockPrefixOp>
- __device__ __forceinline__ void ExclusiveScan(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate, ///< [out] Threadblock-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a threadblock-wide prefix to be applied to all inputs.
- /// Computes an exclusive threadblock-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- __device__ __forceinline__ void ExclusiveSum(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- T &block_aggregate) ///< [out] Threadblock-wide aggregate reduction of input items
- /// Computes an exclusive threadblock-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Instead of using 0 as the threadblock-wide prefix, the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- template <typename BlockPrefixOp>
- __device__ __forceinline__ void ExclusiveSum(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- T &block_aggregate, ///< [out] Threadblock-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a threadblock-wide prefix to be applied to all inputs.
- /// Computes an inclusive threadblock-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- template <typename ScanOp>
- __device__ __forceinline__ void InclusiveScan(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate) ///< [out] Threadblock-wide aggregate reduction of input items
- /// Computes an inclusive threadblock-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- template <
- typename ScanOp,
- typename BlockPrefixOp>
- __device__ __forceinline__ void InclusiveScan(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate, ///< [out] Threadblock-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a threadblock-wide prefix to be applied to all inputs.
- /// Computes an inclusive threadblock-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- __device__ __forceinline__ void InclusiveSum(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- T &block_aggregate) ///< [out] Threadblock-wide aggregate reduction of input items
- /// Computes an inclusive threadblock-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. Instead of using 0 as the threadblock-wide prefix, the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- template <typename BlockPrefixOp>
- __device__ __forceinline__ void InclusiveSum(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- T &block_aggregate, ///< [out] Threadblock-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a threadblock-wide prefix to be applied to all inputs.
- /// Update the calling thread's partial reduction with the warp-wide aggregates from preceding warps. Also returns block-wide aggregate in <em>thread</em><sub>0</sub>.
- /// Computes an exclusive threadblock-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- template <typename ScanOp>
- __device__ __forceinline__ void ExclusiveScan(
- T input, ///< [in] Calling thread's input items
- T &output, ///< [out] Calling thread's output items (may be aliased to \p input)
- const T &identity, ///< [in] Identity value
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate) ///< [out] Threadblock-wide aggregate reduction of input items
- /// Computes an exclusive threadblock-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- template <
- typename ScanOp,
- typename BlockPrefixOp>
- __device__ __forceinline__ void ExclusiveScan(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- T identity, ///< [in] Identity value
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate, ///< [out] Threadblock-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a threadblock-wide prefix to be applied to all inputs.
- /// Computes an exclusive threadblock-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. Also provides every thread with the block-wide \p block_aggregate of all inputs. With no identity value, the output computed for <em>thread</em><sub>0</sub> is undefined.
- template <typename ScanOp>
- __device__ __forceinline__ void ExclusiveScan(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate) ///< [out] Threadblock-wide aggregate reduction of input items
- /// Computes an exclusive threadblock-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- template <
- typename ScanOp,
- typename BlockPrefixOp>
- __device__ __forceinline__ void ExclusiveScan(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate, ///< [out] Threadblock-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a threadblock-wide prefix to be applied to all inputs.
- /// Computes an exclusive threadblock-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- __device__ __forceinline__ void ExclusiveSum(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- T &block_aggregate) ///< [out] Threadblock-wide aggregate reduction of input items
- /// Computes an exclusive threadblock-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Instead of using 0 as the threadblock-wide prefix, the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- template <typename BlockPrefixOp>
- __device__ __forceinline__ void ExclusiveSum(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- T &block_aggregate, ///< [out] Threadblock-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a threadblock-wide prefix to be applied to all inputs.
- /// Computes an inclusive threadblock-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- template <typename ScanOp>
- __device__ __forceinline__ void InclusiveScan(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate) ///< [out] Threadblock-wide aggregate reduction of input items
- /// Computes an inclusive threadblock-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- template <
- typename ScanOp,
- typename BlockPrefixOp>
- __device__ __forceinline__ void InclusiveScan(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T &block_aggregate, ///< [out] Threadblock-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a threadblock-wide prefix to be applied to all inputs.
- /// Computes an inclusive threadblock-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- __device__ __forceinline__ void InclusiveSum(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- T &block_aggregate) ///< [out] Threadblock-wide aggregate reduction of input items
- /// Computes an inclusive threadblock-wide prefix scan using the specified binary \p scan_op functor. Each thread contributes one input element. Instead of using 0 as the threadblock-wide prefix, the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide \p block_aggregate of all inputs.
- template <typename BlockPrefixOp>
- __device__ __forceinline__ void InclusiveSum(
- T input, ///< [in] Calling thread's input item
- T &output, ///< [out] Calling thread's output item (may be aliased to \p input)
- T &block_aggregate, ///< [out] Threadblock-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
- BlockPrefixOp &block_prefix_op) ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a threadblock-wide prefix to be applied to all inputs.
- * cub::BlockHistogramTiles implements a stateful abstraction of CUDA thread blocks for histogramming multiple tiles as part of device-wide histogram.
- int BINS, ///< Number of histogram bins per channel
- int CHANNELS, ///< Number of channels interleaved in the input data (may be greater than the number of active channels being histogrammed)
- int ACTIVE_CHANNELS, ///< Number of channels actively being histogrammed
- typename InputIteratorRA, ///< The input iterator type (may be a simple pointer type). Must have a value type that can be cast as an integer in the range [0..BINS-1]
- typename HistoCounter, ///< Integral type for counting sample occurrences per histogram bin
- RADIX_SORT_SCATTER_DIRECT, ///< Scatter directly from registers to global bins
- RADIX_SORT_SCATTER_TWO_PHASE, ///< First scatter from registers into shared memory bins, then into global bins
-};
-
-
-/**
- * Tuning policy for BlockRadixSortDownsweepTiles
- */
-template <
- int _BLOCK_THREADS, ///< The number of threads per CTA
- int _ITEMS_PER_THREAD, ///< The number of consecutive downsweep keys to process per thread
- BlockLoadAlgorithm _LOAD_ALGORITHM, ///< The BlockLoad algorithm to use
- PtxLoadModifier _LOAD_MODIFIER, ///< The PTX cache-modifier to use for loads
- bool _EXCHANGE_TIME_SLICING, ///< Whether or not to time-slice key/value exchanges through shared memory to lower shared memory pressure
- bool _MEMOIZE_OUTER_SCAN, ///< Whether or not to buffer outer raking scan partials to incur fewer shared memory reads at the expense of higher register pressure. See BlockScanAlgorithm::BLOCK_SCAN_RAKING_MEMOIZE for more details.
- BlockScanAlgorithm _INNER_SCAN_ALGORITHM, ///< The cub::BlockScanAlgorithm algorithm to use
- RadixSortScatterAlgorithm _SCATTER_ALGORITHM, ///< The scattering strategy to use
- * cub::BlockHistogramTilesGlobalAtomic implements a stateful abstraction of CUDA thread blocks for histogramming multiple tiles as part of device-wide histogram.
- */
-
-#pragma once
-
-#include <iterator>
-
-#include "../../../util_type.cuh"
-#include "../../../util_namespace.cuh"
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-
-
-/**
- * BlockHistogramTilesGlobalAtomic implements a stateful abstraction of CUDA thread blocks for histogramming multiple tiles as part of device-wide histogram using global atomics
- int BINS, ///< Number of histogram bins per channel
- int CHANNELS, ///< Number of channels interleaved in the input data (may be greater than the number of active channels being histogrammed)
- int ACTIVE_CHANNELS, ///< Number of channels actively being histogrammed
- typename InputIteratorRA, ///< The input iterator type (may be a simple pointer type). Must have a value type that can be cast as an integer in the range [0..BINS-1]
- typename HistoCounter, ///< Integral type for counting sample occurrences per histogram bin
- * cub::BlockHistogramTilesSharedAtomic implements a stateful abstraction of CUDA thread blocks for histogramming multiple tiles as part of device-wide histogram using shared atomics
- */
-
-#pragma once
-
-#include <iterator>
-
-#include "../../../util_type.cuh"
-#include "../../../util_namespace.cuh"
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-
-/**
- * BlockHistogramTilesSharedAtomic implements a stateful abstraction of CUDA thread blocks for histogramming multiple tiles as part of device-wide histogram using shared atomics
- int CHANNELS, ///< Number of channels interleaved in the input data (may be greater than the number of active channels being histogrammed)
- int ACTIVE_CHANNELS, ///< Number of channels actively being histogrammed
- typename InputIteratorRA, ///< The input iterator type (may be a simple pointer type). Must have a value type that can be cast as an integer in the range [0..BINS-1]
- typename HistoCounter, ///< Integral type for counting sample occurrences per histogram bin
- /// Shared memory type required by this thread block
- struct _TempStorage
- {
- HistoCounter histograms[ACTIVE_CHANNELS][BINS + 1]; // One word of padding between channel histograms to prevent warps working on different histograms from hammering on the same bank
- };
-
-
- /// Alias wrapper allowing storage to be unioned
- * cub::BlockHistogramTilesSort implements a stateful abstraction of CUDA thread blocks for histogramming multiple tiles as part of device-wide histogram using local sorting
- * BlockHistogramTilesSort implements a stateful abstraction of CUDA thread blocks for histogramming multiple tiles as part of device-wide histogram using local sorting
- int BINS, ///< Number of histogram bins per channel
- int CHANNELS, ///< Number of channels interleaved in the input data (may be greater than the number of active channels being histogrammed)
- int ACTIVE_CHANNELS, ///< Number of channels actively being histogrammed
- typename InputIteratorRA, ///< The input iterator type (may be a simple pointer type). Must have a value type that can be cast as an integer in the range [0..BINS-1]
- typename HistoCounter, ///< Integral type for counting sample occurrences per histogram bin
- if (threadIdx.x == 0) grid_queue.ResetDrain(num_samples);
-}
-
-
-/**
- * Histogram pass kernel entry point (multi-block). Computes privatized histograms, one per thread block.
- */
-template <
- typename BlockHistogramTilesPolicy, ///< Tuning policy for cub::BlockHistogramTiles abstraction
- int BINS, ///< Number of histogram bins per channel
- int CHANNELS, ///< Number of channels interleaved in the input data (may be greater than the number of channels being actively histogrammed)
- int ACTIVE_CHANNELS, ///< Number of channels actively being histogrammed
- typename InputIteratorRA, ///< The input iterator type (may be a simple pointer type). Must have a value type that is assignable to <tt>unsigned char</tt>
- typename HistoCounter, ///< Integral type for counting sample occurrences per histogram bin
- typename SizeT> ///< Integer type used for global array indexing
- InputIteratorRA d_samples, ///< [in] Array of sample data. The samples from different channels are assumed to be interleaved (e.g., an array of 32b pixels where each pixel consists of four RGBA 8b samples).
- ArrayWrapper<HistoCounter*, ACTIVE_CHANNELS> d_out_histograms, ///< [out] Histogram counter data having logical dimensions <tt>HistoCounter[ACTIVE_CHANNELS][gridDim.x][BINS]</tt>
- SizeT num_samples, ///< [in] Total number of samples \p d_samples for all channels
- GridEvenShare<SizeT> even_share, ///< [in] Descriptor for how to map an even-share of tiles across thread blocks
- GridQueue<SizeT> queue) ///< [in] Descriptor for performing dynamic mapping of tile data to thread blocks
- * \brief DeviceHistogram provides device-wide parallel operations for constructing histogram(s) from samples data residing within global memory. 
- * \ingroup DeviceModule
- *
- * \par Overview
- * A <a href="http://en.wikipedia.org/wiki/Histogram"><em>histogram</em></a>
- * counts the number of observations that fall into each of the disjoint categories (known as <em>bins</em>).
- *
- * \par Usage Considerations
- * \cdp_class{DeviceHistogram}
- *
- * \par Performance
- *
- * \image html histo_perf.png
- *
- */
-struct DeviceHistogram
-{
-#ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
- * Internal dispatch routine for invoking device-wide, multi-channel, histogram
- */
- template <
- int BINS, ///< Number of histogram bins per channel
- int CHANNELS, ///< Number of channels interleaved in the input data (may be greater than the number of channels being actively histogrammed)
- int ACTIVE_CHANNELS, ///< Number of channels actively being histogrammed
- typename InitHistoKernelPtr, ///< Function type of cub::InitHistoKernel
- typename MultiBlockHistogramKernelPtr, ///< Function type of cub::MultiBlockHistogramKernel
- typename AggregateHistoKernelPtr, ///< Function type of cub::AggregateHistoKernel
- typename InputIteratorRA, ///< The input iterator type (may be a simple pointer type). Must have a value type that is assignable to <tt>unsigned char</tt>
- typename HistoCounter, ///< Integral type for counting sample occurrences per histogram bin
- typename SizeT> ///< Integer type used for global array indexing
- __host__ __device__ __forceinline__
- static cudaError_t Dispatch(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- InitHistoKernelPtr init_kernel, ///< [in] Kernel function pointer to parameterization of cub::InitHistoKernel
- MultiBlockHistogramKernelPtr multi_block_kernel, ///< [in] Kernel function pointer to parameterization of cub::MultiBlockHistogramKernel
- AggregateHistoKernelPtr aggregate_kernel, ///< [in] Kernel function pointer to parameterization of cub::AggregateHistoKernel
- KernelDispachParams &multi_block_dispatch_params, ///< [in] Dispatch parameters that match the policy that \p multi_block_kernel was compiled for
- InputIteratorRA d_samples, ///< [in] Input samples to histogram
- HistoCounter *d_histograms[ACTIVE_CHANNELS], ///< [out] Array of channel histograms, each having BINS counters of integral type \p HistoCounter.
- SizeT num_samples, ///< [in] Number of samples to process
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Default is \p false.
- {
-#ifndef CUB_RUNTIME_ENABLED
-
- // Kernel launch not supported from this device
- return CubDebug(cudaErrorNotSupported);
-
-#else
-
- cudaError error = cudaSuccess;
- do
- {
- // Get device ordinal
- int device_ordinal;
- if (CubDebug(error = cudaGetDevice(&device_ordinal))) break;
-
- // Get SM count
- int sm_count;
- if (CubDebug(error = cudaDeviceGetAttribute (&sm_count, cudaDevAttrMultiProcessorCount, device_ordinal))) break;
-
- // Get a rough estimate of multi_block_kernel SM occupancy based upon the maximum SM occupancy of the targeted PTX architecture
- if (stream_synchronous && (CubDebug(error = SyncStream(stream)))) break;
- }
- }
- while (0);
-
- return error;
-#endif // CUB_RUNTIME_ENABLED
- }
-
-
- /**
- * \brief Computes a device-wide histogram
- *
- * \tparam GRID_ALGORITHM cub::BlockHistogramTilesAlgorithm enumerator specifying the underlying algorithm to use
- * \tparam CHANNELS Number of channels interleaved in the input data (may be greater than the number of channels being actively histogrammed)
- * \tparam ACTIVE_CHANNELS <b>[inferred]</b> Number of channels actively being histogrammed
- * \tparam InputIteratorRA <b>[inferred]</b> Random-access iterator type for input (may be a simple pointer type) Must have a value type that is assignable to <tt>unsigned char</tt>
- * \tparam HistoCounter <b>[inferred]</b> Integral type for counting sample occurrences per histogram bin
- */
- template <
- BlockHistogramTilesAlgorithm GRID_ALGORITHM,
- int BINS, ///< Number of histogram bins per channel
- int CHANNELS, ///< Number of channels interleaved in the input data (may be greater than the number of channels being actively histogrammed)
- int ACTIVE_CHANNELS, ///< Number of channels actively being histogrammed
- typename InputIteratorRA, ///< The input iterator type (may be a simple pointer type). Must have a value type that is assignable to <tt>unsigned char</tt>
- typename HistoCounter> ///< Integral type for counting sample occurrences per histogram bin
- __host__ __device__ __forceinline__
- static cudaError_t Dispatch(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- InputIteratorRA d_samples, ///< [in] Input samples to histogram
- HistoCounter *d_histograms[ACTIVE_CHANNELS], ///< [out] Array of channel histograms, each having BINS counters of integral type \p HistoCounter.
- int num_samples, ///< [in] Number of samples to process
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Default is \p false.
- {
- // Type used for array indexing
- typedef int SizeT;
-
- // Tuning polices for the PTX architecture that will get dispatched to
- * \brief Computes a device-wide histogram. Uses fast block-sorting to compute the histogram. Delivers consistent throughput regardless of sample diversity, but occupancy may be limited by histogram bin count.
- *
- * However, because histograms are privatized in shared memory, a large
- * number of bins (e.g., thousands) may adversely affect occupancy and
- * performance (or even the ability to launch).
- *
- * \devicestorage
- *
- * \cdp
- *
- * \iterator
- *
- * \par
- * The code snippet below illustrates the computation of a 256-bin histogram of
- * \tparam BINS Number of histogram bins per channel
- * \tparam InputIteratorRA <b>[inferred]</b> Random-access iterator type for input (may be a simple pointer type) Must have a value type that can be cast as an integer in the range [0..BINS-1]
- * \tparam HistoCounter <b>[inferred]</b> Integral type for counting sample occurrences per histogram bin
- */
- template <
- int BINS,
- typename InputIteratorRA,
- typename HistoCounter>
- __host__ __device__ __forceinline__
- static cudaError_t SingleChannelSorting(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- HistoCounter* d_histogram, ///< [out] Array of BINS counters of integral type \p HistoCounter.
- int num_samples, ///< [in] Number of samples to process
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Default is \p false.
- * \brief Computes a device-wide histogram. Uses shared-memory atomic read-modify-write operations to compute the histogram. Input samples having lower diversity can cause performance to be degraded, and occupancy may be limited by histogram bin count.
- *
- * However, because histograms are privatized in shared memory, a large
- * number of bins (e.g., thousands) may adversely affect occupancy and
- * performance (or even the ability to launch).
- *
- * \devicestorage
- *
- * \cdp
- *
- * \iterator
- *
- * \par
- * The code snippet below illustrates the computation of a 256-bin histogram of
- * \tparam BINS Number of histogram bins per channel
- * \tparam InputIteratorRA <b>[inferred]</b> Random-access iterator type for input (may be a simple pointer type) Must have a value type that can be cast as an integer in the range [0..BINS-1]
- * \tparam HistoCounter <b>[inferred]</b> Integral type for counting sample occurrences per histogram bin
- */
- template <
- int BINS,
- typename InputIteratorRA,
- typename HistoCounter>
- __host__ __device__ __forceinline__
- static cudaError_t SingleChannelSharedAtomic(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- HistoCounter* d_histogram, ///< [out] Array of BINS counters of integral type \p HistoCounter.
- int num_samples, ///< [in] Number of samples to process
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_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.
- * \brief Computes a device-wide histogram. Uses global-memory atomic read-modify-write operations to compute the histogram. Input samples having lower diversity can cause performance to be degraded.
- *
- * Performance is not significantly impacted when computing histograms having large
- * numbers of bins (e.g., thousands).
- *
- * \devicestorage
- *
- * \cdp
- *
- * \iterator
- *
- * \par
- * The code snippet below illustrates the computation of a 256-bin histogram of
- * \tparam BINS Number of histogram bins per channel
- * \tparam InputIteratorRA <b>[inferred]</b> Random-access iterator type for input (may be a simple pointer type) Must have a value type that can be cast as an integer in the range [0..BINS-1]
- * \tparam HistoCounter <b>[inferred]</b> Integral type for counting sample occurrences per histogram bin
- */
- template <
- int BINS,
- typename InputIteratorRA,
- typename HistoCounter>
- __host__ __device__ __forceinline__
- static cudaError_t SingleChannelGlobalAtomic(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- HistoCounter* d_histogram, ///< [out] Array of BINS counters of integral type \p HistoCounter.
- int num_samples, ///< [in] Number of samples to process
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_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.
- * \brief Computes a device-wide histogram from multi-channel data. Uses fast block-sorting to compute the histogram. Delivers consistent throughput regardless of sample diversity, but occupancy may be limited by histogram bin count.
- *
- * However, because histograms are privatized in shared memory, a large
- * number of bins (e.g., thousands) may adversely affect occupancy and
- * performance (or even the ability to launch).
- *
- * The total number of samples across all channels (\p num_samples) must be a whole multiple of \p CHANNELS.
- *
- * \devicestorage
- *
- * \cdp
- *
- * \iterator
- *
- * \par
- * The code snippet below illustrates the computation of three 256-bin histograms from
- * \tparam BINS Number of histogram bins per channel
- * \tparam CHANNELS Number of channels interleaved in the input data (may be greater than the number of channels being actively histogrammed)
- * \tparam ACTIVE_CHANNELS <b>[inferred]</b> Number of channels actively being histogrammed
- * \tparam InputIteratorRA <b>[inferred]</b> Random-access iterator type for input (may be a simple pointer type) Must have a value type that can be cast as an integer in the range [0..BINS-1]
- * \tparam HistoCounter <b>[inferred]</b> Integral type for counting sample occurrences per histogram bin
- */
- template <
- int BINS,
- int CHANNELS,
- int ACTIVE_CHANNELS,
- typename InputIteratorRA,
- typename HistoCounter>
- __host__ __device__ __forceinline__
- static cudaError_t MultiChannelSorting(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- InputIteratorRA d_samples, ///< [in] Input samples. The samples from different channels are assumed to be interleaved (e.g., an array of 32b pixels where each pixel consists of four RGBA 8b samples).
- HistoCounter *d_histograms[ACTIVE_CHANNELS], ///< [out] Array of channel histogram counter arrays, each having BINS counters of integral type \p HistoCounter.
- int num_samples, ///< [in] Total number of samples to process in all channels, including non-active channels
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_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.
- * \brief Computes a device-wide histogram from multi-channel data. Uses shared-memory atomic read-modify-write operations to compute the histogram. Input samples having lower diversity can cause performance to be degraded, and occupancy may be limited by histogram bin count.
- *
- * However, because histograms are privatized in shared memory, a large
- * number of bins (e.g., thousands) may adversely affect occupancy and
- * performance (or even the ability to launch).
- *
- * The total number of samples across all channels (\p num_samples) must be a whole multiple of \p CHANNELS.
- *
- * \devicestorage
- *
- * \cdp
- *
- * \iterator
- *
- * \par
- * The code snippet below illustrates the computation of three 256-bin histograms from
- * \tparam BINS Number of histogram bins per channel
- * \tparam CHANNELS Number of channels interleaved in the input data (may be greater than the number of channels being actively histogrammed)
- * \tparam ACTIVE_CHANNELS <b>[inferred]</b> Number of channels actively being histogrammed
- * \tparam InputIteratorRA <b>[inferred]</b> Random-access iterator type for input (may be a simple pointer type) Must have a value type that can be cast as an integer in the range [0..BINS-1]
- * \tparam HistoCounter <b>[inferred]</b> Integral type for counting sample occurrences per histogram bin
- */
- template <
- int BINS,
- int CHANNELS,
- int ACTIVE_CHANNELS,
- typename InputIteratorRA,
- typename HistoCounter>
- __host__ __device__ __forceinline__
- static cudaError_t MultiChannelSharedAtomic(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- InputIteratorRA d_samples, ///< [in] Input samples. The samples from different channels are assumed to be interleaved (e.g., an array of 32b pixels where each pixel consists of four RGBA 8b samples).
- HistoCounter *d_histograms[ACTIVE_CHANNELS], ///< [out] Array of channel histogram counter arrays, each having BINS counters of integral type \p HistoCounter.
- int num_samples, ///< [in] Total number of samples to process in all channels, including non-active channels
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_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.
- * \brief Computes a device-wide histogram from multi-channel data. Uses global-memory atomic read-modify-write operations to compute the histogram. Input samples having lower diversity can cause performance to be degraded.
- *
- * Performance is not significantly impacted when computing histograms having large
- * numbers of bins (e.g., thousands).
- *
- * The total number of samples across all channels (\p num_samples) must be a whole multiple of \p CHANNELS.
- *
- * \devicestorage
- *
- * \cdp
- *
- * \iterator
- *
- * Performance is often improved when referencing input samples through a texture-caching iterator, e.g., cub::TexIteratorRA or cub::TexTransformIteratorRA.
- *
- * \par
- * The code snippet below illustrates the computation of three 256-bin histograms from
- * \tparam BINS Number of histogram bins per channel
- * \tparam CHANNELS Number of channels interleaved in the input data (may be greater than the number of channels being actively histogrammed)
- * \tparam ACTIVE_CHANNELS <b>[inferred]</b> Number of channels actively being histogrammed
- * \tparam InputIteratorRA <b>[inferred]</b> Random-access iterator type for input (may be a simple pointer type) Must have a value type that can be cast as an integer in the range [0..BINS-1]
- * \tparam HistoCounter <b>[inferred]</b> Integral type for counting sample occurrences per histogram bin
- */
- template <
- int BINS,
- int CHANNELS,
- int ACTIVE_CHANNELS,
- typename InputIteratorRA,
- typename HistoCounter>
- __host__ __device__ __forceinline__
- static cudaError_t MultiChannelGlobalAtomic(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- InputIteratorRA d_samples, ///< [in] Input samples. The samples from different channels are assumed to be interleaved (e.g., an array of 32b pixels where each pixel consists of four RGBA 8b samples).
- HistoCounter *d_histograms[ACTIVE_CHANNELS], ///< [out] Array of channel histogram counter arrays, each having BINS counters of integral type \p HistoCounter.
- int num_samples, ///< [in] Total number of samples to process in all channels, including non-active channels
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_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.
- SizeT *d_spine, ///< [in,out] Privatized (per block) digit histograms (striped, i.e., 0s counts from each block, then 1s counts from each block, etc.)
- int num_counts) ///< [in] Total number of bin-counts
-{
- // Parameterize the BlockScanTiles type for the current configuration
- Value *d_values_in, ///< [in] Input values ping buffer
- Value *d_values_out, ///< [in] Output values pong buffer
- SizeT *d_spine, ///< [in] Scan of privatized (per block) digit histograms (striped, i.e., 0s counts from each block, then 1s counts from each block, etc.)
- SizeT num_items, ///< [in] Total number of input data items
- int current_bit, ///< [in] Bit position of current radix digit
- bool use_primary_bit_granularity, ///< [in] Whether nor not to use the primary policy (or the embedded alternate policy for smaller bit granularity)
- bool first_pass, ///< [in] Whether this is the first digit pass
- bool last_pass, ///< [in] Whether this is the last digit pass
- GridEvenShare<SizeT> even_share) ///< [in] Descriptor for how to map an even-share of tiles across thread blocks
- * \brief DeviceRadixSort provides operations for computing a device-wide, parallel radix sort across data items residing within global memory. 
- * \ingroup DeviceModule
- *
- * \par Overview
- * The [<em>radix sorting method</em>](http://en.wikipedia.org/wiki/Radix_sort) arranges
- * items into ascending order. It relies upon a positional representation for
- * keys, i.e., each key is comprised of an ordered sequence of symbols (e.g., digits,
- * characters, etc.) specified from least-significant to most-significant. For a
- * given input sequence of keys and a set of rules specifying a total ordering
- * of the symbolic alphabet, the radix sorting method produces a lexicographic
- * ordering of those keys.
- *
- * \par
- * DeviceRadixSort can sort all of the built-in C++ numeric primitive types, e.g.:
- * <tt>unsigned char</tt>, \p int, \p double, etc. Although the direct radix sorting
- * method can only be applied to unsigned integral types, BlockRadixSort
- * is able to sort signed and floating-point types via simple bit-wise transformations
- * that ensure lexicographic key ordering.
- *
- * \par Usage Considerations
- * \cdp_class{DeviceRadixSort}
- *
- * \par Performance
- *
- * \image html lsd_sort_perf.png
- *
- */
-struct DeviceRadixSort
-{
- #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
- * Internal dispatch routine for computing a device-wide reduction using a two-stages of kernel invocations.
- */
- template <
- typename UpsweepKernelPtr, ///< Function type of cub::RadixSortUpsweepKernel
- typename SpineKernelPtr, ///< Function type of cub::SpineScanKernel
- typename DownsweepKernelPtr, ///< Function type of cub::RadixSortUpsweepKernel
- typename Key, ///< Key type
- typename Value, ///< Value type
- typename SizeT> ///< Integer type used for global array indexing
- __host__ __device__ __forceinline__
- static cudaError_t Dispatch(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- UpsweepKernelPtr upsweep_kernel, ///< [in] Kernel function pointer to parameterization of cub::RadixSortUpsweepKernel
- SpineKernelPtr scan_kernel, ///< [in] Kernel function pointer to parameterization of cub::SpineScanKernel
- DownsweepKernelPtr downsweep_kernel, ///< [in] Kernel function pointer to parameterization of cub::RadixSortUpsweepKernel
- KernelDispachParams &upsweep_dispatch_params, ///< [in] Dispatch parameters that match the policy that \p upsweep_kernel was compiled for
- KernelDispachParams &scan_dispatch_params, ///< [in] Dispatch parameters that match the policy that \p scan_kernel was compiled for
- KernelDispachParams &downsweep_dispatch_params, ///< [in] Dispatch parameters that match the policy that \p downsweep_kernel was compiled for
- DoubleBuffer<Key> &d_keys, ///< [in,out] Double-buffer whose current buffer contains the unsorted input keys and, upon return, is updated to point to the sorted output keys
- DoubleBuffer<Value> &d_values, ///< [in,out] Double-buffer whose current buffer contains the unsorted input values and, upon return, is updated to point to the sorted output values
- SizeT num_items, ///< [in] Number of items to reduce
- int begin_bit = 0, ///< [in] <b>[optional]</b> The beginning (least-significant) bit index needed for key comparison
- int end_bit = sizeof(Key) * 8, ///< [in] <b>[optional]</b> The past-the-end (most-significant) bit index needed for key comparison
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Default is \p false.
- {
-#ifndef CUB_RUNTIME_ENABLED
-
- // Kernel launch not supported from this device
- return CubDebug(cudaErrorNotSupported );
-
-#else
-
- cudaError error = cudaSuccess;
- do
- {
- // Get device ordinal
- int device_ordinal;
- if (CubDebug(error = cudaGetDevice(&device_ordinal))) break;
-
- // Get SM count
- int sm_count;
- if (CubDebug(error = cudaDeviceGetAttribute (&sm_count, cudaDevAttrMultiProcessorCount, device_ordinal))) break;
-
- // Get a rough estimate of downsweep_kernel SM occupancy based upon the maximum SM occupancy of the targeted PTX architecture
- * // Sorted keys and values are referenced by d_keys.Current() and d_values.Current()
- *
- * \endcode
- *
- * \tparam Key <b>[inferred]</b> Key type
- * \tparam Value <b>[inferred]</b> Value type
- */
- template <
- typename Key,
- typename Value>
- __host__ __device__ __forceinline__
- static cudaError_t SortPairs(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- DoubleBuffer<Key> &d_keys, ///< [in,out] Double-buffer of keys whose current buffer contains the unsorted input keys and, upon return, is updated to point to the sorted output keys
- DoubleBuffer<Value> &d_values, ///< [in,out] Double-buffer of values whose current buffer contains the unsorted input values and, upon return, is updated to point to the sorted output values
- int num_items, ///< [in] Number of items to reduce
- int begin_bit = 0, ///< [in] <b>[optional]</b> The first (least-significant) bit index needed for key comparison
- int end_bit = sizeof(Key) * 8, ///< [in] <b>[optional]</b> The past-the-end (most-significant) bit index needed for key comparison
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Default is \p false.
- * // Sorted keys are referenced by d_keys.Current()
- *
- * \endcode
- *
- * \tparam Key <b>[inferred]</b> Key type
- */
- template <typename Key>
- __host__ __device__ __forceinline__
- static cudaError_t SortKeys(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- DoubleBuffer<Key> &d_keys, ///< [in,out] Double-buffer of keys whose current buffer contains the unsorted input keys and, upon return, is updated to point to the sorted output keys
- int num_items, ///< [in] Number of items to reduce
- int begin_bit = 0, ///< [in] <b>[optional]</b> The first (least-significant) bit index needed for key comparison
- int end_bit = sizeof(Key) * 8, ///< [in] <b>[optional]</b> The past-the-end (most-significant) bit index needed for key comparison
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Default is \p false.
- * \brief DeviceReduce provides operations for computing a device-wide, parallel reduction across data items residing within global memory. 
- * \ingroup DeviceModule
- *
- * \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 list of input elements.
- *
- * \par Usage Considerations
- * \cdp_class{DeviceReduce}
- *
- * \par Performance
- *
- * \image html reduction_perf.png
- *
- */
-struct DeviceReduce
-{
-#ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
- * Internal dispatch routine for computing a device-wide reduction using a two-stages of kernel invocations.
- */
- template <
- typename ReducePrivatizedKernelPtr, ///< Function type of cub::ReducePrivatizedKernel
- typename ReduceSingleKernelPtr, ///< Function type of cub::ReduceSingleKernel
- typename ResetDrainKernelPtr, ///< Function type of cub::ResetDrainKernel
- typename InputIteratorRA, ///< Random-access iterator type for input (may be a simple pointer type)
- typename OutputIteratorRA, ///< Random-access iterator type for output (may be a simple pointer type)
- typename SizeT, ///< Integer type used for global array indexing
- typename ReductionOp> ///< Binary reduction operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- __host__ __device__ __forceinline__
- static cudaError_t Dispatch(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- ReducePrivatizedKernelPtr privatized_kernel, ///< [in] Kernel function pointer to parameterization of cub::ReducePrivatizedKernel
- ReduceSingleKernelPtr single_kernel, ///< [in] Kernel function pointer to parameterization of cub::ReduceSingleKernel
- ResetDrainKernelPtr prepare_drain_kernel, ///< [in] Kernel function pointer to parameterization of cub::ResetDrainKernel
- KernelDispachParams &privatized_dispatch_params, ///< [in] Dispatch parameters that match the policy that \p privatized_kernel_ptr was compiled for
- KernelDispachParams &single_dispatch_params, ///< [in] Dispatch parameters that match the policy that \p single_kernel was compiled for
- InputIteratorRA d_in, ///< [in] Input data to reduce
- OutputIteratorRA d_out, ///< [out] Output location for result
- SizeT num_items, ///< [in] Number of items to reduce
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Default is \p false.
- * \tparam InputIteratorRA <b>[inferred]</b> Random-access iterator type for input (may be a simple pointer type)
- * \tparam OutputIteratorRA <b>[inferred]</b> Random-access iterator type for output (may be a simple pointer type)
- * \tparam ReductionOp <b>[inferred]</b> Binary reduction operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- */
- template <
- typename InputIteratorRA,
- typename OutputIteratorRA,
- typename ReductionOp>
- __host__ __device__ __forceinline__
- static cudaError_t Reduce(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- InputIteratorRA d_in, ///< [in] Input data to reduce
- OutputIteratorRA d_out, ///< [out] Output location for result
- int num_items, ///< [in] Number of items to reduce
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Default is \p false.
- * \tparam InputIteratorRA <b>[inferred]</b> Random-access iterator type for input (may be a simple pointer type)
- * \tparam OutputIteratorRA <b>[inferred]</b> Random-access iterator type for output (may be a simple pointer type)
- */
- template <
- typename InputIteratorRA,
- typename OutputIteratorRA>
- __host__ __device__ __forceinline__
- static cudaError_t Sum(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- InputIteratorRA d_in, ///< [in] Input data to reduce
- OutputIteratorRA d_out, ///< [out] Output location for result
- int num_items, ///< [in] Number of items to reduce
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Default is \p false.
- * \brief DeviceReduceByKey provides operations for computing a device-wide, parallel prefix scan across data items residing within global memory. 
- */
-struct DeviceReduceByKey
-{
-#ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
- typename InitScanKernelPtr, ///< Function type of cub::InitScanKernel
- typename MultiBlockScanKernelPtr, ///< Function type of cub::MultiBlockScanKernel
- typename InputIteratorRA, ///< Random-access iterator type for input (may be a simple pointer type)
- typename OutputIteratorRA, ///< Random-access iterator type for output (may be a simple pointer type)
- typename ReductionOp, ///< Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- typename Identity, ///< Identity value type (cub::NullType for inclusive scans)
- typename SizeT> ///< Integer type used for global array indexing
- __host__ __device__ __forceinline__
- static cudaError_t Dispatch(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- InitScanKernelPtr init_kernel, ///< [in] Kernel function pointer to parameterization of cub::InitScanKernel
- MultiBlockScanKernelPtr multi_block_kernel, ///< [in] Kernel function pointer to parameterization of cub::MultiBlockScanKernel
- KernelDispachParams &multi_block_dispatch_params, ///< [in] Dispatch parameters that match the policy that \p multi_block_kernel was compiled for
- InputIteratorRA d_in, ///< [in] Iterator pointing to scan input
- OutputIteratorRA d_out, ///< [in] Iterator pointing to scan output
- SizeT num_items, ///< [in] Total number of items to scan
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Default is \p false.
- if (stream_synchronous && CubDebug(error = cudaStreamSynchronize(stream))) break;
-#else
- if (stream_synchronous && CubDebug(error = cudaDeviceSynchronize())) break;
-#endif
- }
- while (0);
-
- return error;
-
-#endif // CUB_RUNTIME_ENABLED
- }
-
-
-
- /**
- * Internal scan dispatch routine for using default tuning policies
- */
- template <
- typename InputIteratorRA, ///< Random-access iterator type for input (may be a simple pointer type)
- typename OutputIteratorRA, ///< Random-access iterator type for output (may be a simple pointer type)
- typename ReductionOp, ///< Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- typename Identity, ///< Identity value type (cub::NullType for inclusive scans)
- typename SizeT> ///< Integer type used for global array indexing
- __host__ __device__ __forceinline__
- static cudaError_t Dispatch(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- InputIteratorRA d_in, ///< [in] Iterator pointing to scan input
- OutputIteratorRA d_out, ///< [in] Iterator pointing to scan output
- SizeT num_items, ///< [in] Total number of items to scan
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Default is \p false.
- * \brief Computes device-wide reductions of consecutive values whose corresponding keys are equal.
- *
- * The resulting output lists of value-aggregates and their corresponding keys are compacted.
- *
- * \devicestorage
- *
- * \tparam KeyInputIteratorRA <b>[inferred]</b> Random-access input iterator type for keys input (may be a simple pointer type)
- * \tparam KeyOutputIteratorRA <b>[inferred]</b> Random-access output iterator type for keys output (may be a simple pointer type)
- * \tparam ValueInputIteratorRA <b>[inferred]</b> Random-access input iterator type for values input (may be a simple pointer type)
- * \tparam ValueOutputIteratorRA <b>[inferred]</b> Random-access output iterator type for values output (may be a simple pointer type)
- * \tparam ReductionOp <b>[inferred]</b> Binary reduction operator type having member <tt>T operator()(const T &a, const T &b)</tt>, where \p T is the value type of \p ValueInputIteratorRA
- */
- template <
- typename KeyInputIteratorRA,
- typename KeyOutputIteratorRA,
- typename ValueInputIteratorRA,
- typename ValueOutputIteratorRA,
- typename ReductionOp>
- __host__ __device__ __forceinline__
- static cudaError_t ReduceValues(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- KeyInputIteratorRA d_keys_in, ///< [in] Key input data
- KeyOutputIteratorRA d_keys_out, ///< [out] Key output data (compacted)
- ValueInputIteratorRA d_values_in, ///< [in] Value input data
- ValueOutputIteratorRA d_values_out, ///< [out] Value output data (compacted)
- int num_items, ///< [in] Total number of input pairs
- ReductionOp reduction_op, ///< [in] Binary value reduction operator
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_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.
- * \brief Computes device-wide sums of consecutive values whose corresponding keys are equal.
- *
- * The resulting output lists of value-aggregates and their corresponding keys are compacted.
- *
- * \devicestorage
- *
- * \tparam KeyInputIteratorRA <b>[inferred]</b> Random-access input iterator type for keys input (may be a simple pointer type)
- * \tparam KeyOutputIteratorRA <b>[inferred]</b> Random-access output iterator type for keys output (may be a simple pointer type)
- * \tparam ValueInputIteratorRA <b>[inferred]</b> Random-access input iterator type for values input (may be a simple pointer type)
- * \tparam ValueOutputIteratorRA <b>[inferred]</b> Random-access output iterator type for values output (may be a simple pointer type)
- * \tparam ReductionOp <b>[inferred]</b> Binary reduction operator type having member <tt>T operator()(const T &a, const T &b)</tt>, where \p T is the value type of \p ValueInputIteratorRA
- */
- template <
- typename KeyInputIteratorRA,
- typename KeyOutputIteratorRA,
- typename ValueInputIteratorRA,
- typename ValueOutputIteratorRA>
- __host__ __device__ __forceinline__
- static cudaError_t SumValues(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- KeyInputIteratorRA d_keys_in, ///< [in] Key input data
- KeyOutputIteratorRA d_keys_out, ///< [in] Key output data (compacted)
- ValueInputIteratorRA d_values_in, ///< [in] Value input data
- ValueOutputIteratorRA d_values_out, ///< [in] Value output data (compacted)
- int num_items, ///< [in] Total number of input pairs
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_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.
- * \brief Computes the "run-length" of each group of consecutive, equal-valued keys.
- *
- * The resulting output lists of run-length counts and their corresponding keys are compacted.
- *
- * \devicestorage
- *
- * \tparam KeyInputIteratorRA <b>[inferred]</b> Random-access input iterator type for keys input (may be a simple pointer type)
- * \tparam KeyOutputIteratorRA <b>[inferred]</b> Random-access output iterator type for keys output (may be a simple pointer type)
- * \tparam CountOutputIteratorRA <b>[inferred]</b> Random-access output iterator type for output of key-counts whose value type must be convertible to an integer type (may be a simple pointer type)
- */
- template <
- typename KeyInputIteratorRA,
- typename KeyOutputIteratorRA,
- typename CountOutputIteratorRA>
- __host__ __device__ __forceinline__
- static cudaError_t RunLengths(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- KeyInputIteratorRA d_keys_in, ///< [in] Key input data
- KeyOutputIteratorRA d_keys_out, ///< [in] Key output data (compacted)
- CountOutputIteratorRA d_counts_out, ///< [in] Run-length counts output data (compacted)
- int num_items, ///< [in] Total number of keys
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_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.
- * \brief Removes duplicates within each group of consecutive, equal-valued keys. Only the first key from each group (and corresponding value) is kept.
- *
- * The resulting keys are compacted.
- *
- * \devicestorage
- *
- * \tparam KeyInputIteratorRA <b>[inferred]</b> Random-access input iterator type for keys input (may be a simple pointer type)
- * \tparam KeyOutputIteratorRA <b>[inferred]</b> Random-access output iterator type for keys output (may be a simple pointer type)
- * \tparam ValueInputIteratorRA <b>[inferred]</b> Random-access input iterator type for values input (may be a simple pointer type)
- * \tparam ValueOutputIteratorRA <b>[inferred]</b> Random-access output iterator type for values output (may be a simple pointer type)
- * \tparam ReductionOp <b>[inferred]</b> Binary reduction operator type having member <tt>T operator()(const T &a, const T &b)</tt>, where \p T is the value type of \p ValueInputIteratorRA
- */
- template <
- typename KeyInputIteratorRA,
- typename KeyOutputIteratorRA,
- typename ValueInputIteratorRA,
- typename ValueOutputIteratorRA,
- typename ReductionOp>
- __host__ __device__ __forceinline__
- static cudaError_t Unique(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- KeyInputIteratorRA d_keys_in, ///< [in] Key input data
- KeyOutputIteratorRA d_keys_out, ///< [out] Key output data (compacted)
- ValueInputIteratorRA d_values_in, ///< [in] Value input data
- ValueOutputIteratorRA d_values_out, ///< [out] Value output data (compacted)
- int num_items, ///< [in] Total number of input pairs
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_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.
-#ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
-
-/**
- * Partition kernel entry point (multi-block)
- */
-template <
- typename BlockPartitionTilesPolicy, ///< Tuning policy for cub::BlockPartitionTiles abstraction
- typename InputIteratorRA, ///< Random-access iterator type for input (may be a simple pointer type)
- typename OutputIteratorRA, ///< Random-access iterator type for output (may be a simple pointer type)
- typename LengthOutputIterator, ///< Output iterator type for recording the length of the first partition (may be a simple pointer type)
- typename PredicateOp, ///< Unary predicate operator indicating membership in the first partition type having member <tt>bool operator()(const T &val)</tt>
- typename SizeT> ///< Integer type used for global array indexing
- typename ScanInitKernelPtr, ///< Function type of cub::ScanInitKernel
- typename PartitionKernelPtr, ///< Function type of cub::PartitionKernel
- typename InputIteratorRA, ///< Random-access iterator type for input (may be a simple pointer type)
- typename OutputIteratorRA, ///< Random-access iterator type for output (may be a simple pointer type)
- typename LengthOutputIterator, ///< Output iterator type for recording the length of the first partition (may be a simple pointer type)
- typename PredicateOp, ///< Unary predicate operator indicating membership in the first partition type having member <tt>bool operator()(const T &val)</tt>
- typename SizeT> ///< Integer type used for global array indexing
- __host__ __device__ __forceinline__
- static cudaError_t Dispatch(
- int ptx_version, ///< [in] PTX version
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- ScanInitKernelPtr init_kernel, ///< [in] Kernel function pointer to parameterization of cub::PartitionInitKernel
- PartitionKernelPtr partition_kernel, ///< [in] Kernel function pointer to parameterization of cub::PartitionKernel
- KernelDispachParams &scan_dispatch_params, ///< [in] Dispatch parameters that match the policy that \p partition_kernel was compiled for
- InputIteratorRA d_in, ///< [in] Iterator pointing to scan input
- OutputIteratorRA d_out, ///< [in] Iterator pointing to scan output
- LengthOutputIterator d_partition_length, ///< [out] Output iterator referencing the location where the pivot offset (i.e., the length of the first partition) is to be recorded
- PredicateOp pred_op, ///< [in] Unary predicate operator indicating membership in the first partition
- SizeT num_items, ///< [in] Total number of items to partition
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Default is \p false.
- if (stream_synchronous && (CubDebug(error = SyncStream(stream)))) break;
- }
- while (0);
-
- return error;
-
-#endif // CUB_RUNTIME_ENABLED
- }
-
-
-
- /**
- * Internal partition dispatch routine for using default tuning policies
- */
- template <
- typename PARTITIONS, ///< Number of partitions we are keeping
- typename InputIteratorRA, ///< Random-access iterator type for input (may be a simple pointer type)
- typename OutputIteratorRA, ///< Random-access iterator type for output (may be a simple pointer type)
- typename LengthOutputIterator, ///< Output iterator type for recording the length of the first partition (may be a simple pointer type)
- typename PredicateOp, ///< Unary predicate operator indicating membership in the first partition type having member <tt>bool operator()(const T &val)</tt>
- typename SizeT> ///< Integer type used for global array indexing
- __host__ __device__ __forceinline__
- static cudaError_t Dispatch(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- InputIteratorRA d_in, ///< [in] Iterator pointing to input items
- OutputIteratorRA d_out, ///< [in] Iterator pointing to output items
- LengthOutputIterator d_partition_length, ///< [out] Output iterator referencing the location where the pivot offset (i.e., the length of the first partition) is to be recorded
- PredicateOp pred_op, ///< [in] Unary predicate operator indicating membership in the first partition
- SizeT num_items, ///< [in] Total number of items to partition
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Default is \p false.
- * \brief Splits a list of input items into two partitions within the given output list using the specified predicate. The relative ordering of inputs is not necessarily preserved.
- *
- * An item \p val is placed in the first partition if <tt>pred_op(val) == true</tt>, otherwise
- * it is placed in the second partition. The offset of the partitioning pivot (equivalent to
- * the total length of the first partition as well as the starting offset of the second), is
- * recorded to \p d_partition_length.
- *
- * The length of the output referenced by \p d_out is assumed to be the same as that of \p d_in.
- *
- * \devicestorage
- *
- * \tparam InputIteratorRA <b>[inferred]</b> Random-access iterator type for input (may be a simple pointer type)
- * \tparam OutputIteratorRA <b>[inferred]</b> Random-access iterator type for output (may be a simple pointer type)
- * \tparam LengthOutputIterator <b>[inferred]</b> Random-access iterator type for output (may be a simple pointer type)
- * \tparam PredicateOp <b>[inferred]</b> Unary predicate operator indicating membership in the first partition type having member <tt>bool operator()(const T &val)</tt>
- */
- template <
- typename InputIteratorRA,
- typename OutputIteratorRA,
- typename LengthOutputIterator,
- typename PredicateOp>
- __host__ __device__ __forceinline__
- static cudaError_t Partition(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- InputIteratorRA d_in, ///< [in] Iterator pointing to input items
- OutputIteratorRA d_out, ///< [in] Iterator pointing to output items
- LengthOutputIterator d_pivot_offset, ///< [out] Output iterator referencing the location where the pivot offset is to be recorded
- PredicateOp pred_op, ///< [in] Unary predicate operator indicating membership in the first partition
- int num_items, ///< [in] Total number of items to partition
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_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.
- * \brief DeviceScan provides operations for computing a device-wide, parallel prefix scan across data items residing within global memory. 
- * \ingroup DeviceModule
- *
- * \par Overview
- * Given a list of input elements and a binary reduction operator, a [<em>prefix scan</em>](http://en.wikipedia.org/wiki/Prefix_sum)
- * produces an output list where each element is computed to be the reduction
- * of the elements occurring earlier in the input list. <em>Prefix sum</em>
- * connotes a prefix scan with the addition operator. The term \em inclusive indicates
- * that the <em>i</em><sup>th</sup> output reduction incorporates the <em>i</em><sup>th</sup> input.
- * The term \em exclusive indicates the <em>i</em><sup>th</sup> input is not incorporated into
- * the <em>i</em><sup>th</sup> output reduction.
- *
- * \par Usage Considerations
- * \cdp_class{DeviceScan}
- *
- * \par Performance
- *
- * \image html scan_perf.png
- *
- */
-struct DeviceScan
-{
-#ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
- typename ScanInitKernelPtr, ///< Function type of cub::ScanInitKernel
- typename ScanKernelPtr, ///< Function type of cub::ScanKernel
- typename InputIteratorRA, ///< Random-access iterator type for input (may be a simple pointer type)
- typename OutputIteratorRA, ///< Random-access iterator type for output (may be a simple pointer type)
- typename ScanOp, ///< Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- typename Identity, ///< Identity value type (cub::NullType for inclusive scans)
- typename SizeT> ///< Integer type used for global array indexing
- __host__ __device__ __forceinline__
- static cudaError_t Dispatch(
- int ptx_version, ///< [in] PTX version
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- ScanInitKernelPtr init_kernel, ///< [in] Kernel function pointer to parameterization of cub::ScanInitKernel
- ScanKernelPtr scan_kernel, ///< [in] Kernel function pointer to parameterization of cub::ScanKernel
- KernelDispachParams &scan_dispatch_params, ///< [in] Dispatch parameters that match the policy that \p scan_kernel was compiled for
- InputIteratorRA d_in, ///< [in] Iterator pointing to scan input
- OutputIteratorRA d_out, ///< [in] Iterator pointing to scan output
- ScanOp scan_op, ///< [in] Binary scan operator
- Identity identity, ///< [in] Identity element
- SizeT num_items, ///< [in] Total number of items to scan
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Default is \p false.
- if (stream_synchronous && (CubDebug(error = SyncStream(stream)))) break;
- }
- while (0);
-
- return error;
-
-#endif // CUB_RUNTIME_ENABLED
- }
-
-
-
- /**
- * Internal scan dispatch routine for using default tuning policies
- */
- template <
- typename InputIteratorRA, ///< Random-access iterator type for input (may be a simple pointer type)
- typename OutputIteratorRA, ///< Random-access iterator type for output (may be a simple pointer type)
- typename ScanOp, ///< Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- typename Identity, ///< Identity value type (cub::NullType for inclusive scans)
- typename SizeT> ///< Integer type used for global array indexing
- __host__ __device__ __forceinline__
- static cudaError_t Dispatch(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- InputIteratorRA d_in, ///< [in] Iterator pointing to scan input
- OutputIteratorRA d_out, ///< [in] Iterator pointing to scan output
- ScanOp scan_op, ///< [in] Binary scan operator
- Identity identity, ///< [in] Identity element
- SizeT num_items, ///< [in] Total number of items to scan
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_synchronous = false) ///< [in] <b>[optional]</b> Whether or not to synchronize the stream after every kernel launch to check for errors. Default is \p false.
- * \tparam InputIteratorRA <b>[inferred]</b> Random-access iterator type for input (may be a simple pointer type)
- * \tparam OutputIteratorRA <b>[inferred]</b> Random-access iterator type for output (may be a simple pointer type)
- */
- template <
- typename InputIteratorRA,
- typename OutputIteratorRA>
- __host__ __device__ __forceinline__
- static cudaError_t ExclusiveSum(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- InputIteratorRA d_in, ///< [in] Iterator pointing to scan input
- OutputIteratorRA d_out, ///< [in] Iterator pointing to scan output
- int num_items, ///< [in] Total number of items to scan
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_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.
- * \tparam InputIteratorRA <b>[inferred]</b> Random-access iterator type for input (may be a simple pointer type)
- * \tparam OutputIteratorRA <b>[inferred]</b> Random-access iterator type for output (may be a simple pointer type)
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- * \tparam Identity <b>[inferred]</b> Type of the \p identity value used Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- */
- template <
- typename InputIteratorRA,
- typename OutputIteratorRA,
- typename ScanOp,
- typename Identity>
- __host__ __device__ __forceinline__
- static cudaError_t ExclusiveScan(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- InputIteratorRA d_in, ///< [in] Iterator pointing to scan input
- OutputIteratorRA d_out, ///< [in] Iterator pointing to scan output
- ScanOp scan_op, ///< [in] Binary scan operator
- Identity identity, ///< [in] Identity element
- int num_items, ///< [in] Total number of items to scan
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_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.
- * \tparam InputIteratorRA <b>[inferred]</b> Random-access iterator type for input (may be a simple pointer type)
- * \tparam OutputIteratorRA <b>[inferred]</b> Random-access iterator type for output (may be a simple pointer type)
- */
- template <
- typename InputIteratorRA,
- typename OutputIteratorRA>
- __host__ __device__ __forceinline__
- static cudaError_t InclusiveSum(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- InputIteratorRA d_in, ///< [in] Iterator pointing to scan input
- OutputIteratorRA d_out, ///< [in] Iterator pointing to scan output
- int num_items, ///< [in] Total number of items to scan
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_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.
- * \tparam InputIteratorRA <b>[inferred]</b> Random-access iterator type for input (may be a simple pointer type)
- * \tparam OutputIteratorRA <b>[inferred]</b> Random-access iterator type for output (may be a simple pointer type)
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- */
- template <
- typename InputIteratorRA,
- typename OutputIteratorRA,
- typename ScanOp>
- __host__ __device__ __forceinline__
- static cudaError_t InclusiveScan(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \p d_temp_storage allocation.
- InputIteratorRA d_in, ///< [in] Iterator pointing to scan input
- OutputIteratorRA d_out, ///< [in] Iterator pointing to scan output
- ScanOp scan_op, ///< [in] Binary scan operator
- int num_items, ///< [in] Total number of items to scan
- cudaStream_t stream = 0, ///< [in] <b>[optional]</b> CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
- bool stream_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.
- * cub::GridEvenShare is a descriptor utility for distributing input among CUDA threadblocks in an "even-share" fashion. Each threadblock gets roughly the same number of fixed-size work units (grains).
- */
-
-
-#pragma once
-
-#include "../util_namespace.cuh"
-#include "../util_macro.cuh"
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-
-/**
- * \addtogroup GridModule
- * @{
- */
-
-
-/**
- * \brief GridEvenShare is a descriptor utility for distributing input among CUDA threadblocks in an "even-share" fashion. Each threadblock gets roughly the same number of fixed-size work units (grains).
- *
- * \par Overview
- * GridEvenShare indicates which sections of input are to be mapped onto which threadblocks.
- * Threadblocks may receive one of three different amounts of work: "big", "normal",
- * and "last". The "big" workloads are one scheduling grain larger than "normal". The "last" work unit
- * for the last threadblock may be partially-full if the input is not an even multiple of
- * the scheduling grain size.
- *
- * \par
- * Before invoking a child grid, a parent thread will typically construct and initialize an instance of
- * GridEvenShare using \p GridInit(). The instance can be passed to child threadblocks which can
- * initialize their per-threadblock offsets using \p BlockInit().
- *
- * \tparam SizeT Integer type for array indexing
- */
-template <typename SizeT>
-class GridEvenShare
-{
-private:
-
- SizeT total_grains;
- int big_blocks;
- SizeT big_share;
- SizeT normal_share;
- SizeT normal_base_offset;
-
-
-public:
-
- /// Total number of input items
- SizeT num_items;
-
- /// Grid size in threadblocks
- int grid_size;
-
- /// Offset into input marking the beginning of the owning thread block's segment of input tiles
- SizeT block_offset;
-
- /// Offset into input of marking the end (one-past) of the owning thread block's segment of input tiles
- SizeT block_oob;
-
- /**
- * \brief Block-based constructor for single-block grids.
- SizeT num_items, ///< Total number of input items
- int max_grid_size, ///< Maximum grid size allowable (actual grid size may be less if not warranted by the the number of input items)
- int schedule_granularity) ///< Granularity by which the input can be parcelled into and distributed among threablocks. Usually the thread block's native tile size (or a multiple thereof.
- * \brief cub::GridMappingStrategy enumerates alternative strategies for mapping constant-sized tiles of device-wide data onto a grid of CUDA thread blocks.
- */
-enum GridMappingStrategy
-{
- /**
- * \brief An "even-share" strategy for assigning input tiles to thread blocks.
- *
- * \par Overview
- * The input is evenly partitioned into \p p segments, where \p p is
- * constant and corresponds loosely to the number of thread blocks that may
- * actively reside on the target device. Each segment is comprised of
- * consecutive tiles, where a tile is a small, constant-sized unit of input
- * to be processed to completion before the thread block terminates or
- * obtains more work. The kernel invokes \p p thread blocks, each
- * of which iteratively consumes a segment of <em>n</em>/<em>p</em> elements
- * in tile-size increments.
- */
- GRID_MAPPING_EVEN_SHARE,
-
- /**
- * \brief A dynamic "queue-based" strategy for assigning input tiles to thread blocks.
- *
- * \par Overview
- * The input is treated as a queue to be dynamically consumed by a grid of
- * thread blocks. Work is atomically dequeued in tiles, where a tile is a
- * unit of input to be processed to completion before the thread block
- * terminates or obtains more work. The grid size \p p is constant,
- * loosely corresponding to the number of thread blocks that may actively
- * cub::GridQueue is a descriptor utility for dynamic queue management.
- */
-
-#pragma once
-
-#include "../util_namespace.cuh"
-#include "../util_debug.cuh"
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-
-/**
- * \addtogroup GridModule
- * @{
- */
-
-
-/**
- * \brief GridQueue is a descriptor utility for dynamic queue management.
- *
- * \par Overview
- * GridQueue descriptors provides abstractions for "filling" or
- * "draining" globally-shared vectors.
- *
- * \par
- * A "filling" GridQueue works by atomically-adding to a zero-initialized counter,
- * returning a unique offset for the calling thread to write its items.
- * The GridQueue maintains the total "fill-size". The fill counter must be reset
- * using GridQueue::ResetFill by the host or kernel instance prior to the kernel instance that
- * will be filling.
- *
- * \par
- * Similarly a "draining" GridQueue works by works by atomically-incrementing a
- * zero-initialized counter, returning a unique offset for the calling thread to
- * read its items. Threads can safely drain until the array's logical fill-size is
- * exceeded. The drain counter must be reset using GridQueue::ResetDrain or
- * GridQueue::ResetDrainAfterFill by the host or kernel instance prior to the kernel instance that
- * will be filling. (For dynamic work distribution of existing data, the corresponding fill-size
- * is simply the number of elements in the array.)
- *
- * \par
- * Iterative work management can be implemented simply with a pair of flip-flopping
- * work buffers, each with an associated set of fill and drain GridQueue descriptors.
- *
- * \tparam SizeT Integer type for array indexing
- */
-template <typename SizeT>
-class GridQueue
-{
-private:
-
- /// Counter indices
- enum
- {
- FILL = 0,
- DRAIN = 1,
- };
-
- /// Pair of counters
- SizeT *d_counters;
-
-public:
-
- /// Returns the device allocation size in bytes needed to construct a GridQueue instance
- __host__ __device__ __forceinline__
- static size_t AllocationSize()
- {
- return sizeof(SizeT) * 2;
- }
-
-
- /// Constructs an invalid GridQueue descriptor around the device storage allocation
- __host__ __device__ __forceinline__ GridQueue(
- void *d_storage) ///< Device allocation to back the GridQueue. Must be at least as big as <tt>AllocationSize()</tt>.
- :
- d_counters((SizeT*) d_storage)
- {}
-
-
- /// This operation resets the drain so that it may advance to meet the existing fill-size. To be called by the host or by a kernel prior to that which will be draining.
- /// This operation sets the fill-size and resets the drain counter, preparing the GridQueue for draining in the next kernel instance. To be called by the host or by a kernel prior to that which will be draining.
- * Thread utilities for sequential reduction over statically-sized array types
- */
-
-#pragma once
-
-#include "../thread/thread_operators.cuh"
-#include "../util_namespace.cuh"
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-/**
- * \addtogroup ThreadModule
- * @{
- */
-
-/**
- * \name Sequential reduction over statically-sized array types
- * @{
- */
-
-/**
- * \brief Perform a sequential reduction over \p LENGTH elements of the \p input array, seeded with the specified \p prefix. The aggregate is returned.
- *
- * \tparam LENGTH Length of input array
- * \tparam T <b>[inferred]</b> The data type to be reduced.
- * \tparam ScanOp <b>[inferred]</b> Binary reduction operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- * Thread utilities for sequential prefix scan over statically-sized array types
- */
-
-#pragma once
-
-#include "../thread/thread_operators.cuh"
-#include "../util_namespace.cuh"
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-/**
- * \addtogroup ThreadModule
- * @{
- */
-
-/**
- * \name Sequential prefix scan over statically-sized array types
- * @{
- */
-
-/**
- * \brief Perform a sequential exclusive prefix scan over \p LENGTH elements of the \p input array, seeded with the specified \p prefix. The aggregate is returned.
- *
- * \tparam LENGTH Length of \p input and \p output arrays
- * \tparam T <b>[inferred]</b> The data type to be scanned.
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- */
-template <
- int LENGTH,
- typename T,
- typename ScanOp>
-__device__ __forceinline__ T ThreadScanExclusive(
- T *input, ///< [in] Input array
- T *output, ///< [out] Output array (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T prefix, ///< [in] Prefix to seed scan with
- bool apply_prefix = true) ///< [in] Whether or not the calling thread should apply its prefix. If not, the first output element is undefined. (Handy for preventing thread-0 from applying a prefix.)
-{
- T inclusive = input[0];
- if (apply_prefix)
- {
- inclusive = scan_op(prefix, inclusive);
- }
- output[0] = prefix;
- T exclusive = inclusive;
-
- #pragma unroll
- for (int i = 1; i < LENGTH; ++i)
- {
- inclusive = scan_op(exclusive, input[i]);
- output[i] = exclusive;
- exclusive = inclusive;
- }
-
- return inclusive;
-}
-
-
-/**
- * \brief Perform a sequential exclusive prefix scan over the statically-sized \p input array, seeded with the specified \p prefix. The aggregate is returned.
- *
- * \tparam LENGTH <b>[inferred]</b> Length of \p input and \p output arrays
- * \tparam T <b>[inferred]</b> The data type to be scanned.
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- */
-template <
- int LENGTH,
- typename T,
- typename ScanOp>
-__device__ __forceinline__ T ThreadScanExclusive(
- T (&input)[LENGTH], ///< [in] Input array
- T (&output)[LENGTH], ///< [out] Output array (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T prefix, ///< [in] Prefix to seed scan with
- bool apply_prefix = true) ///< [in] Whether or not the calling thread should apply its prefix. (Handy for preventing thread-0 from applying a prefix.)
- * \brief Perform a sequential inclusive prefix scan over \p LENGTH elements of the \p input array, seeded with the specified \p prefix. The aggregate is returned.
- *
- * \tparam LENGTH Length of \p input and \p output arrays
- * \tparam T <b>[inferred]</b> The data type to be scanned.
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- */
-template <
- int LENGTH,
- typename T,
- typename ScanOp>
-__device__ __forceinline__ T ThreadScanInclusive(
- T *input, ///< [in] Input array
- T *output, ///< [out] Output array (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T prefix, ///< [in] Prefix to seed scan with
- bool apply_prefix = true) ///< [in] Whether or not the calling thread should apply its prefix. (Handy for preventing thread-0 from applying a prefix.)
-{
- T inclusive = input[0];
- if (apply_prefix)
- {
- inclusive = scan_op(prefix, inclusive);
- }
- output[0] = inclusive;
-
- // Continue scan
- #pragma unroll
- for (int i = 1; i < LENGTH; ++i)
- {
- inclusive = scan_op(inclusive, input[i]);
- output[i] = inclusive;
- }
-
- return inclusive;
-}
-
-
-/**
- * \brief Perform a sequential inclusive prefix scan over the statically-sized \p input array, seeded with the specified \p prefix. The aggregate is returned.
- *
- * \tparam LENGTH <b>[inferred]</b> Length of \p input and \p output arrays
- * \tparam T <b>[inferred]</b> The data type to be scanned.
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- */
-template <
- int LENGTH,
- typename T,
- typename ScanOp>
-__device__ __forceinline__ T ThreadScanInclusive(
- T (&input)[LENGTH], ///< [in] Input array
- T (&output)[LENGTH], ///< [out] Output array (may be aliased to \p input)
- ScanOp scan_op, ///< [in] Binary scan operator
- T prefix, ///< [in] Prefix to seed scan with
- bool apply_prefix = true) ///< [in] Whether or not the calling thread should apply its prefix. (Handy for preventing thread-0 from applying a prefix.)
- Spinlock spin_lock; /// Spinlock for thread-safety
-
- unsigned int bin_growth; /// Geometric growth factor for bin-sizes
- unsigned int min_bin; /// Minimum bin enumeration
- unsigned int max_bin; /// Maximum bin enumeration
-
- size_t min_bin_bytes; /// Minimum bin size
- size_t max_bin_bytes; /// Maximum bin size
- size_t max_cached_bytes; /// Maximum aggregate cached bytes per device
-
- bool debug; /// Whether or not to print (de)allocation events to stdout
- bool skip_cleanup; /// Whether or not to skip a call to FreeAllCached() when destructor is called. (The CUDA runtime may have already shut down for statically declared allocators)
-
-#ifndef __CUDA_ARCH__ // Only define STL container members in host code
-
- GpuCachedBytes cached_bytes; /// Map of device ordinal to aggregate cached bytes on that device
- CachedBlocks cached_blocks; /// Set of cached device allocations available for reuse
- BusyBlocks live_blocks; /// Set of live device allocations currently in use
- * The following macros definitions are supported:
- * - \p CUB_LOG. Simple event messages are printed to \p stdout.
- */
-
-#pragma once
-
-#include <stdio.h>
-#include "util_namespace.cuh"
-#include "util_arch.cuh"
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-
-/**
- * \addtogroup UtilModule
- * @{
- */
-
-
-/// CUB error reporting macro (prints error messages to stderr)
-#if (defined(DEBUG) || defined(_DEBUG))
- #define CUB_STDERR
-#endif
-
-
-
-/**
- * \brief %If \p CUB_STDERR is defined and \p error is not \p cudaSuccess, the corresponding error message is printed to \p stderr (or \p stdout in device code) along with the supplied source context.
- * Properties of a given CUDA device and the corresponding PTX bundle
- */
-
-#pragma once
-
-#include "util_arch.cuh"
-#include "util_debug.cuh"
-#include "util_namespace.cuh"
-#include "util_macro.cuh"
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-
-/**
- * \addtogroup UtilModule
- * @{
- */
-
-#ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
-
-
-/**
- * Empty kernel for querying PTX manifest metadata (e.g., version) for the current device
- */
-template <typename T>
-__global__ void EmptyKernel(void) { }
-
-
-/**
- * Alias temporaries to externally-allocated device storage (or simply return the amount of storage needed).
- */
-template <int ALLOCATIONS>
-__host__ __device__ __forceinline__
-cudaError_t AliasTemporaries(
- void *d_temp_storage, ///< [in] %Device allocation of temporary storage. When NULL, the required allocation size is returned in \p temp_storage_bytes and no work is done.
- size_t &temp_storage_bytes, ///< [in,out] Size in bytes of \t d_temp_storage allocation
- void* (&allocations)[ALLOCATIONS], ///< [in,out] Pointers to device allocations needed
- size_t (&allocation_sizes)[ALLOCATIONS]) ///< [in] Sizes in bytes of device allocations needed
-{
- const int ALIGN_BYTES = 256;
- const int ALIGN_MASK = ~(ALIGN_BYTES - 1);
-
- // Compute exclusive prefix sum over allocation requests
- * \brief A simple random-access transform iterator for applying a transformation operator.
- *
- * \par Overview
- * TransformIteratorRA is a random-access iterator that wraps both a native
- * device pointer of type <tt>InputType*</tt> and a unary conversion functor of
- * type \p ConversionOp. \p OutputType references are made by pulling \p InputType
- * values through the \p ConversionOp instance.
- *
- * \tparam InputType The value type of the pointer being wrapped
- * \tparam ConversionOp Unary functor type for mapping objects of type \p InputType to type \p OutputType. Must have member <tt>OutputType operator()(const InputType &datum)</tt>.
- * \tparam OutputType The value type of this iterator
- * \brief A simple random-access iterator for loading primitive values through texture cache.
- *
- * \par Overview
- * TexIteratorRA is a random-access iterator that wraps a native
- * device pointer of type <tt>T*</tt>. References made through TexIteratorRA
- * causes values to be pulled through texture cache.
- *
- * \par Usage Considerations
- * - Can only be used with primitive types (e.g., \p char, \p int, \p float), with the exception of \p double
- * - Only one TexIteratorRA or TexIteratorRA of a certain \p InputType can be bound at any given time (per host thread)
- *
- * \tparam InputType The value type of the pointer being wrapped
- * \tparam ConversionOp Unary functor type for mapping objects of type \p InputType to type \p OutputType. Must have member <tt>OutputType operator()(const InputType &datum)</tt>.
- * \tparam OutputType The value type of this iterator
- */
-template <typename T>
-class TexIteratorRA
-{
-public:
-#ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
- * \brief A simple random-access transform iterator for loading primitive values through texture cache and and subsequently applying a transformation operator.
- *
- * \par Overview
- * TexTransformIteratorRA is a random-access iterator that wraps both a native
- * device pointer of type <tt>InputType*</tt> and a unary conversion functor of
- * type \p ConversionOp. \p OutputType references are made by pulling \p InputType
- * values through the texture cache and then transformed them using the
- * \p ConversionOp instance.
- *
- * \par Usage Considerations
- * - Can only be used with primitive types (e.g., \p char, \p int, \p float), with the exception of \p double
- * - Only one TexIteratorRA or TexTransformIteratorRA of a certain \p InputType can be bound at any given time (per host thread)
- *
- * \tparam InputType The value type of the pointer being wrapped
- * \tparam ConversionOp Unary functor type for mapping objects of type \p InputType to type \p OutputType. Must have member <tt>OutputType operator()(const InputType &datum)</tt>.
- * \tparam OutputType The value type of this iterator
- * \brief Allows for the treatment of an integral constant as a type at compile-time (e.g., to achieve static call dispatch based on constant integral values)
- * \brief A wrapper for passing simple static arrays as kernel parameters
- */
-template <typename T, int COUNT>
-struct ArrayWrapper
-{
- /// Static array of type \p T
- T array[COUNT];
-};
-
-
-/**
- * \brief Double-buffer storage wrapper for multi-pass stream transformations that require more than one storage array for streaming intermediate results back and forth.
- *
- * Many multi-pass computations require a pair of "ping-pong" storage
- * buffers (e.g., one for reading from and the other for writing to, and then
- * vice-versa for the subsequent pass). This structure wraps a set of device
- * buffers and a "selector" member to track which is "current".
- */
-template <typename T>
-struct DoubleBuffer
-{
- /// Pair of device buffer pointers
- T *d_buffers[2];
-
- /// Selector into \p d_buffers (i.e., the active/valid buffer)
- * \brief Defines a structure \p detector_name that is templated on type \p T. The \p detector_name struct exposes a constant member \p VALUE indicating whether or not parameter \p T exposes a nested type \p nested_type_name
- * \brief Determine whether or not BinaryOp's functor is of the form <tt>bool operator()(const T& a, const T&b)</tt> or <tt>bool operator()(const T& a, const T&b, unsigned int idx)</tt>
- */
-template <typename T, typename BinaryOp>
-struct BinaryOpHasIdxParam
-{
-private:
- template <typename BinaryOpT, bool (BinaryOpT::*)(const T &a, const T &b, unsigned int idx) const> struct SFINAE1 {};
- template <typename BinaryOpT, bool (BinaryOpT::*)(const T &a, const T &b, unsigned int idx)> struct SFINAE2 {};
- template <typename BinaryOpT, bool (BinaryOpT::*)(T a, T b, unsigned int idx) const> struct SFINAE3 {};
- template <typename BinaryOpT, bool (BinaryOpT::*)(T a, T b, unsigned int idx)> struct SFINAE4 {};
-
- template <typename BinaryOpT, bool (BinaryOpT::*)(const T &a, const T &b, int idx) const> struct SFINAE5 {};
- template <typename BinaryOpT, bool (BinaryOpT::*)(const T &a, const T &b, int idx)> struct SFINAE6 {};
- template <typename BinaryOpT, bool (BinaryOpT::*)(T a, T b, int idx) const> struct SFINAE7 {};
- template <typename BinaryOpT, bool (BinaryOpT::*)(T a, T b, int idx)> struct SFINAE8 {};
- * \brief Exposes a member typedef \p Type that names the corresponding CUDA vector type if one exists. Otherwise \p Type refers to the VectorHelper structure itself, which will wrap the corresponding \p x, \p y, etc. vector fields.
- */
-template <typename T, int vec_elements> struct VectorHelper;
-
-#ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
-
-enum
-{
- /// The maximum number of elements in CUDA vector types
- MAX_VEC_ELEMENTS = 4,
-};
-
-
-/**
- * Generic vector-1 type
- */
-template <typename T>
-struct VectorHelper<T, 1>
-{
- enum { BUILT_IN = false };
-
- T x;
-
- typedef VectorHelper<T, 1> Type;
-};
-
-/**
- * Generic vector-2 type
- */
-template <typename T>
-struct VectorHelper<T, 2>
-{
- enum { BUILT_IN = false };
-
- T x;
- T y;
-
- typedef VectorHelper<T, 2> Type;
-};
-
-/**
- * Generic vector-3 type
- */
-template <typename T>
-struct VectorHelper<T, 3>
-{
- enum { BUILT_IN = false };
-
- T x;
- T y;
- T z;
-
- typedef VectorHelper<T, 3> Type;
-};
-
-/**
- * Generic vector-4 type
- */
-template <typename T>
-struct VectorHelper<T, 4>
-{
- enum { BUILT_IN = false };
-
- T x;
- T y;
- T z;
- T w;
-
- typedef VectorHelper<T, 4> Type;
-};
-
-/**
- * Macro for expanding partially-specialized built-in vector types
- * The cub::WarpReduce class provides [<em>collective</em>](index.html#sec0) methods for computing a parallel reduction of items partitioned across CUDA warp threads.
- */
-
-#pragma once
-
-#include "specializations/warp_reduce_shfl.cuh"
-#include "specializations/warp_reduce_smem.cuh"
-#include "../thread/thread_operators.cuh"
-#include "../util_arch.cuh"
-#include "../util_type.cuh"
-#include "../util_namespace.cuh"
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-
-/**
- * \addtogroup WarpModule
- * @{
- */
-
-/**
- * \brief The WarpReduce class provides [<em>collective</em>](index.html#sec0) methods for computing a parallel reduction of items partitioned across CUDA warp threads. 
- *
- * \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 list of input elements.
- *
- * \tparam T The reduction input/output element type
- * \tparam LOGICAL_WARPS <b>[optional]</b> The number of entrant "logical" warps performing concurrent warp reductions. Default is 1.
- * \tparam LOGICAL_WARP_THREADS <b>[optional]</b> The number of threads per "logical" warp (may be less than the number of hardware warp threads). Default is the warp size of the targeted CUDA compute-capability (e.g., 32 threads for SM20).
- *
- * \par Simple Examples
- * \warpcollective{WarpReduce}
- * \par
- * The code snippet below illustrates four concurrent warp sum reductions within a block of
- * 128 threads (one per each of the 32-thread warps).
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize WarpReduce for 4 warps on type int
- #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
-
- /// Internal specialization. Use SHFL-based reduction if (architecture is >= SM30) and ((only one logical warp) or (LOGICAL_WARP_THREADS is a power-of-two))
- * \brief Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Logical warp and lane identifiers are constructed from <tt>threadIdx.x</tt>.
- * \brief Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Logical warp and lane identifiers are constructed from <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ WarpReduce(
- TempStorage &temp_storage) ///< [in] Reference to memory allocation having layout type TempStorage
- * \brief Collective constructor using a private static allocation of shared memory as temporary storage. Threads are identified using the given warp and lane identifiers.
- */
- __device__ __forceinline__ WarpReduce(
- int warp_id, ///< [in] A suitable warp membership identifier
- int lane_id) ///< [in] A lane identifier within the warp
- :
- temp_storage(PrivateStorage()),
- warp_id(warp_id),
- lane_id(lane_id)
- {}
-
-
- /**
- * \brief Collective constructor using the specified memory allocation as temporary storage. Threads are identified using the given warp and lane identifiers.
- */
- __device__ __forceinline__ WarpReduce(
- TempStorage &temp_storage, ///< [in] Reference to memory allocation having layout type TempStorage
- int warp_id, ///< [in] A suitable warp membership identifier
- int lane_id) ///< [in] A lane identifier within the warp
- * \brief Computes a segmented sum in each active warp where segments are defined by head-flags. The sum of each segment is returned to the first lane in that segment (which always includes <em>lane</em><sub>0</sub>).
- *
- * \smemreuse
- *
- * The code snippet below illustrates a head-segmented warp sum
- * reduction within a block of 32 threads (one warp).
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize WarpReduce for a single warp on type int
- * \brief Computes a segmented sum in each active warp where segments are defined by tail-flags. The sum of each segment is returned to the first lane in that segment (which always includes <em>lane</em><sub>0</sub>).
- *
- * \smemreuse
- *
- * The code snippet below illustrates a tail-segmented warp sum
- * reduction within a block of 32 threads (one warp).
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize WarpReduce for a single warp on type int
- * \brief Computes a warp-wide reduction in each active warp using the specified binary reduction functor. The output is valid in warp <em>lane</em><sub>0</sub>.
- *
- * Supports non-commutative reduction operators
- *
- * \smemreuse
- *
- * The code snippet below illustrates four concurrent warp max reductions within a block of
- * 128 threads (one per each of the 32-thread warps).
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize WarpReduce for 4 warps on type int
- * \brief Computes a partially-full warp-wide reduction in each active warp using the specified binary reduction functor. The output is valid in warp <em>lane</em><sub>0</sub>.
- *
- * All threads in each logical warp must agree on the same value for \p valid_items. Otherwise the result is undefined.
- *
- * Supports non-commutative reduction operators
- *
- * \smemreuse
- *
- * The code snippet below illustrates a max reduction within a single, partially-full
- * block of 32 threads (one warp).
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(int *d_data, int valid_items)
- * {
- * // Specialize WarpReduce for a single warp on type int
- * \brief Computes a segmented reduction in each active warp where segments are defined by head-flags. The reduction of each segment is returned to the first lane in that segment (which always includes <em>lane</em><sub>0</sub>).
- *
- * Supports non-commutative reduction operators
- *
- * \smemreuse
- *
- * The code snippet below illustrates a head-segmented warp max
- * reduction within a block of 32 threads (one warp).
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize WarpReduce for a single warp on type int
- * \brief Computes a segmented reduction in each active warp where segments are defined by tail-flags. The reduction of each segment is returned to the first lane in that segment (which always includes <em>lane</em><sub>0</sub>).
- *
- * Supports non-commutative reduction operators
- *
- * \smemreuse
- *
- * The code snippet below illustrates a tail-segmented warp max
- * reduction within a block of 32 threads (one warp).
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize WarpReduce for a single warp on type int
- * The cub::WarpScan class provides [<em>collective</em>](index.html#sec0) methods for computing a parallel prefix scan of items partitioned across CUDA warp threads.
- */
-
-#pragma once
-
-#include "specializations/warp_scan_shfl.cuh"
-#include "specializations/warp_scan_smem.cuh"
-#include "../thread/thread_operators.cuh"
-#include "../util_arch.cuh"
-#include "../util_type.cuh"
-#include "../util_namespace.cuh"
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-/**
- * \addtogroup WarpModule
- * @{
- */
-
-/**
- * \brief The WarpScan class provides [<em>collective</em>](index.html#sec0) methods for computing a parallel prefix scan of items partitioned across CUDA warp threads. 
- *
- * \par Overview
- * Given a list of input elements and a binary reduction operator, a [<em>prefix scan</em>](http://en.wikipedia.org/wiki/Prefix_sum)
- * produces an output list where each element is computed to be the reduction
- * of the elements occurring earlier in the input list. <em>Prefix sum</em>
- * connotes a prefix scan with the addition operator. The term \em inclusive indicates
- * that the <em>i</em><sup>th</sup> output reduction incorporates the <em>i</em><sup>th</sup> input.
- * The term \em exclusive indicates the <em>i</em><sup>th</sup> input is not incorporated into
- * the <em>i</em><sup>th</sup> output reduction.
- *
- * \tparam T The scan input/output element type
- * \tparam LOGICAL_WARPS <b>[optional]</b> The number of "logical" warps performing concurrent warp scans. Default is 1.
- * \tparam LOGICAL_WARP_THREADS <b>[optional]</b> The number of threads per "logical" warp (may be less than the number of hardware warp threads). Default is the warp size associated with the CUDA Compute Capability targeted by the compiler (e.g., 32 threads for SM20).
- *
- * \par Simple Examples
- * \warpcollective{WarpScan}
- * \par
- * The code snippet below illustrates four concurrent warp prefix sums within a block of
- * 128 threads (one per each of the 32-thread warps).
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize WarpScan for 4 warps on type int
- /// Internal specialization. Use SHFL-based reduction if (architecture is >= SM30) and ((only one logical warp) or (LOGICAL_WARP_THREADS is a power-of-two))
- * \brief Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Logical warp and lane identifiers are constructed from <tt>threadIdx.x</tt>.
- * \brief Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Logical warp and lane identifiers are constructed from <tt>threadIdx.x</tt>.
- */
- __device__ __forceinline__ WarpScan(
- TempStorage &temp_storage) ///< [in] Reference to memory allocation having layout type TempStorage
- * \brief Collective constructor using a private static allocation of shared memory as temporary storage. Threads are identified using the given warp and lane identifiers.
- */
- __device__ __forceinline__ WarpScan(
- int warp_id, ///< [in] A suitable warp membership identifier
- int lane_id) ///< [in] A lane identifier within the warp
- :
- temp_storage(PrivateStorage()),
- warp_id(warp_id),
- lane_id(lane_id)
- {}
-
-
- /**
- * \brief Collective constructor using the specified memory allocation as temporary storage. Threads are identified using the given warp and lane identifiers.
- */
- __device__ __forceinline__ WarpScan(
- TempStorage &temp_storage, ///< [in] Reference to memory allocation having layout type TempStorage
- int warp_id, ///< [in] A suitable warp membership identifier
- int lane_id) ///< [in] A lane identifier within the warp
- * \brief Computes an inclusive prefix sum in each logical warp. Instead of using 0 as the warp-wide prefix, the call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs.
- *
- * The \p warp_aggregate is undefined in threads other than <em>warp-lane</em><sub>0</sub>.
- *
- * The \p warp_prefix_op functor must implement a member function <tt>T operator()(T warp_aggregate)</tt>.
- * The functor's input parameter \p warp_aggregate is the same value also returned by the scan operation.
- * The functor will be invoked by the entire warp of threads, however only the return value from
- * <em>lane</em><sub>0</sub> is applied as the threadblock-wide prefix. Can be stateful.
- *
- * \smemreuse
- *
- * The code snippet below illustrates a single thread block of 32 threads (one warp) that progressively
- * computes an inclusive prefix sum over multiple "tiles" of input using a
- * prefix functor to maintain a running total between block-wide scans. Each tile consists
- * of 32 integer items that are partitioned across the warp.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * // A stateful callback functor that maintains a running prefix to be applied
- * Suppose the input \p d_data is <tt>1, 1, 1, 1, 1, 1, 1, 1, ...</tt>.
- * The corresponding output for the first segment will be <tt>1, 2, 3, ..., 32</tt>.
- * The output for the second segment will be <tt>33, 34, 35, ..., 64</tt>. Furthermore,
- * the value \p 32 will be stored in \p warp_aggregate for all threads after each scan.
- *
- * \tparam WarpPrefixOp <b>[inferred]</b> Call-back functor type having member <tt>T operator()(T warp_aggregate)</tt>
- */
- template <typename WarpPrefixOp>
- __device__ __forceinline__ void InclusiveSum(
- T input, ///< [in] Calling thread's input item.
- T &output, ///< [out] Calling thread's output item. May be aliased with \p input.
- T &warp_aggregate, ///< [out] <b>[<em>warp-lane</em><sub>0</sub> only]</b> Warp-wide aggregate reduction of input items, exclusive of the \p warp_prefix_op value
- WarpPrefixOp &warp_prefix_op) ///< [in-out] <b>[<em>warp-lane</em><sub>0</sub> only]</b> Call-back functor for specifying a warp-wide prefix to be applied to all inputs.
- {
- // Compute inclusive warp scan
- InclusiveSum(input, output, warp_aggregate);
-
- // Compute warp-wide prefix from aggregate, then broadcast to other lanes
- /// Computes an exclusive prefix sum in each logical warp.
- __device__ __forceinline__ void ExclusiveSum(T input, T &output, Int2Type<true> is_primitive)
- {
- // Compute exclusive warp scan from inclusive warp scan
- T inclusive;
- InclusiveSum(input, inclusive);
- output = inclusive - input;
- }
-
- /// Computes an exclusive prefix sum in each logical warp. Specialized for non-primitive types.
- __device__ __forceinline__ void ExclusiveSum(T input, T &output, Int2Type<false> is_primitive)
- {
- // Delegate to regular scan for non-primitive types (because we won't be able to use subtraction)
- T identity = T();
- ExclusiveScan(input, output, identity, Sum());
- }
-
- /// Computes an exclusive prefix sum in each logical warp. Also provides every thread with the warp-wide \p warp_aggregate of all inputs.
- __device__ __forceinline__ void ExclusiveSum(T input, T &output, T &warp_aggregate, Int2Type<true> is_primitive)
- {
- // Compute exclusive warp scan from inclusive warp scan
- T inclusive;
- InclusiveSum(input, inclusive, warp_aggregate);
- output = inclusive - input;
- }
-
- /// Computes an exclusive prefix sum in each logical warp. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. Specialized for non-primitive types.
- __device__ __forceinline__ void ExclusiveSum(T input, T &output, T &warp_aggregate, Int2Type<false> is_primitive)
- {
- // Delegate to regular scan for non-primitive types (because we won't be able to use subtraction)
- /// Computes an exclusive prefix sum in each logical warp. Instead of using 0 as the warp-wide prefix, the call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs.
- template <typename WarpPrefixOp>
- __device__ __forceinline__ void ExclusiveSum(T input, T &output, T &warp_aggregate, WarpPrefixOp &warp_prefix_op, Int2Type<true> is_primitive)
- {
- // Compute exclusive warp scan from inclusive warp scan
- /// Computes an exclusive prefix sum in each logical warp. Instead of using 0 as the warp-wide prefix, the call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. Specialized for non-primitive types.
- template <typename WarpPrefixOp>
- __device__ __forceinline__ void ExclusiveSum(T input, T &output, T &warp_aggregate, WarpPrefixOp &warp_prefix_op, Int2Type<false> is_primitive)
- {
- // Delegate to regular scan for non-primitive types (because we won't be able to use subtraction)
- * \brief Computes an exclusive prefix sum in each logical warp. Instead of using 0 as the warp-wide prefix, the call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs.
- *
- * This operation assumes the value of obtained by the <tt>T</tt>'s default
- * constructor (or by zero-initialization if no user-defined default
- * constructor exists) is suitable as the identity value "zero" for
- * addition.
- *
- * The \p warp_prefix_op functor must implement a member function <tt>T operator()(T warp_aggregate)</tt>.
- * The functor's input parameter \p warp_aggregate is the same value also returned by the scan operation.
- * The functor will be invoked by the entire warp of threads, however only the return value from
- * <em>lane</em><sub>0</sub> is applied as the threadblock-wide prefix. Can be stateful.
- *
- * \smemreuse
- *
- * The code snippet below illustrates a single thread block of 32 threads (one warp) that progressively
- * computes an exclusive prefix sum over multiple "tiles" of input using a
- * prefix functor to maintain a running total between block-wide scans. Each tile consists
- * of 32 integer items that are partitioned across the warp.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * // A stateful callback functor that maintains a running prefix to be applied
- * Suppose the input \p d_data is <tt>1, 1, 1, 1, 1, 1, 1, 1, ...</tt>.
- * The corresponding output for the first segment will be <tt>0, 1, 2, ..., 31</tt>.
- * The output for the second segment will be <tt>32, 33, 34, ..., 63</tt>. Furthermore,
- * the value \p 32 will be stored in \p warp_aggregate for all threads after each scan.
- *
- * \tparam WarpPrefixOp <b>[inferred]</b> Call-back functor type having member <tt>T operator()(T warp_aggregate)</tt>
- */
- template <typename WarpPrefixOp>
- __device__ __forceinline__ void ExclusiveSum(
- T input, ///< [in] Calling thread's input item.
- T &output, ///< [out] Calling thread's output item. May be aliased with \p input.
- T &warp_aggregate, ///< [out] <b>[<em>warp-lane</em><sub>0</sub> only]</b> Warp-wide aggregate reduction of input items (exclusive of the \p warp_prefix_op value).
- WarpPrefixOp &warp_prefix_op) ///< [in-out] <b>[<em>warp-lane</em><sub>0</sub> only]</b> Call-back functor for specifying a warp-wide prefix to be applied to all inputs.
- * \brief Computes an inclusive prefix sum using the specified binary scan functor in each logical warp. Also provides every thread with the warp-wide \p warp_aggregate of all inputs.
- *
- * Supports non-commutative scan operators.
- *
- * \smemreuse
- *
- * The code snippet below illustrates four concurrent warp-wide inclusive prefix max scans within a block of
- * 128 threads (one per each of the 32-thread warps).
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize WarpScan for 4 warps on type int
- * \brief Computes an inclusive prefix sum using the specified binary scan functor in each logical warp. The call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs.
- *
- * The \p warp_prefix_op functor must implement a member function <tt>T operator()(T warp_aggregate)</tt>.
- * The functor's input parameter \p warp_aggregate is the same value also returned by the scan operation.
- * The functor will be invoked by the entire warp of threads, however only the return value from
- * <em>lane</em><sub>0</sub> is applied as the threadblock-wide prefix. Can be stateful.
- *
- * Supports non-commutative scan operators.
- *
- * \smemreuse
- *
- * The code snippet below illustrates a single thread block of 32 threads (one warp) that progressively
- * computes an inclusive prefix max scan over multiple "tiles" of input using a
- * prefix functor to maintain a running total between block-wide scans. Each tile consists
- * of 32 integer items that are partitioned across the warp.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * // A stateful callback functor that maintains a running prefix to be applied
- * Suppose the input \p d_data is <tt>0, -1, 2, -3, 4, -5, ...</tt>.
- * The corresponding output for the first segment will be <tt>0, 0, 2, 2, ..., 30, 30</tt>.
- * The output for the second segment will be <tt>32, 32, 34, 34, ..., 62, 62</tt>. Furthermore,
- * \p block_aggregate will be assigned \p 30 in all threads after the first scan, assigned \p 62 after the second
- * scan, etc.
- *
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- * \tparam WarpPrefixOp <b>[inferred]</b> Call-back functor type having member <tt>T operator()(T warp_aggregate)</tt>
- */
- template <
- typename ScanOp,
- typename WarpPrefixOp>
- __device__ __forceinline__ void InclusiveScan(
- T input, ///< [in] Calling thread's input item.
- T &output, ///< [out] Calling thread's output item. May be aliased with \p input.
- ScanOp scan_op, ///< [in] Binary scan operator
- T &warp_aggregate, ///< [out] <b>[<em>warp-lane</em><sub>0</sub> only]</b> Warp-wide aggregate reduction of input items (exclusive of the \p warp_prefix_op value).
- WarpPrefixOp &warp_prefix_op) ///< [in-out] <b>[<em>warp-lane</em><sub>0</sub> only]</b> Call-back functor for specifying a warp-wide prefix to be applied to all inputs.
- * \brief Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. Also provides every thread with the warp-wide \p warp_aggregate of all inputs.
- *
- * Supports non-commutative scan operators.
- *
- * \smemreuse
- *
- * The code snippet below illustrates four concurrent warp-wide exclusive prefix max scans within a block of
- * 128 threads (one per each of the 32-thread warps).
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize WarpScan for 4 warps on type int
- * \brief Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. The call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs.
- *
- * The \p warp_prefix_op functor must implement a member function <tt>T operator()(T warp_aggregate)</tt>.
- * The functor's input parameter \p warp_aggregate is the same value also returned by the scan operation.
- * The functor will be invoked by the entire warp of threads, however only the return value from
- * <em>lane</em><sub>0</sub> is applied as the threadblock-wide prefix. Can be stateful.
- *
- * Supports non-commutative scan operators.
- *
- * \smemreuse
- *
- * The code snippet below illustrates a single thread block of 32 threads (one warp) that progressively
- * computes an exclusive prefix max scan over multiple "tiles" of input using a
- * prefix functor to maintain a running total between block-wide scans. Each tile consists
- * of 32 integer items that are partitioned across the warp.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * // A stateful callback functor that maintains a running prefix to be applied
- * Suppose the input \p d_data is <tt>0, -1, 2, -3, 4, -5, ...</tt>.
- * The corresponding output for the first segment will be <tt>INT_MIN, 0, 0, 2, ..., 28, 30</tt>.
- * The output for the second segment will be <tt>30, 32, 32, 34, ..., 60, 62</tt>. Furthermore,
- * \p block_aggregate will be assigned \p 30 in all threads after the first scan, assigned \p 62 after the second
- * scan, etc.
- *
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- * \tparam WarpPrefixOp <b>[inferred]</b> Call-back functor type having member <tt>T operator()(T warp_aggregate)</tt>
- */
- template <
- typename ScanOp,
- typename WarpPrefixOp>
- __device__ __forceinline__ void ExclusiveScan(
- T input, ///< [in] Calling thread's input item.
- T &output, ///< [out] Calling thread's output item. May be aliased with \p input.
- T identity, ///< [in] Identity value
- ScanOp scan_op, ///< [in] Binary scan operator
- T &warp_aggregate, ///< [out] <b>[<em>warp-lane</em><sub>0</sub> only]</b> Warp-wide aggregate reduction of input items (exclusive of the \p warp_prefix_op value).
- WarpPrefixOp &warp_prefix_op) ///< [in-out] <b>[<em>warp-lane</em><sub>0</sub> only]</b> Call-back functor for specifying a warp-wide prefix to be applied to all inputs.
- * \brief Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. Because no identity value is supplied, the \p output computed for <em>warp-lane</em><sub>0</sub> is undefined.
- *
- * Supports non-commutative scan operators.
- *
- * \smemreuse
- *
- * The code snippet below illustrates four concurrent warp-wide exclusive prefix max scans within a block of
- * 128 threads (one per each of the 32-thread warps).
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize WarpScan for 4 warps on type int
- * \brief Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. Because no identity value is supplied, the \p output computed for <em>warp-lane</em><sub>0</sub> is undefined. Also provides every thread with the warp-wide \p warp_aggregate of all inputs.
- *
- * Supports non-commutative scan operators.
- *
- * \smemreuse
- *
- * The code snippet below illustrates four concurrent warp-wide exclusive prefix max scans within a block of
- * 128 threads (one per each of the 32-thread warps).
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- * // Specialize WarpScan for 4 warps on type int
- * Suppose the set of input \p thread_data across the block of threads is <tt>0, -1, 2, -3, ..., 126, -127</tt>.
- * The corresponding output \p thread_data in the first warp would be
- * <tt>?, 0, 0, 2, ..., 28, 30</tt>, the output for the second warp would be <tt>?, 32, 32, 34, ..., 60, 62</tt>, etc.
- * (The output \p thread_data in each warp lane0 is undefined.) Furthermore, \p warp_aggregate would be assigned \p 30 for threads in the first warp, \p 62 for threads
- * in the second warp, etc.
- *
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- */
- template <typename ScanOp>
- __device__ __forceinline__ void ExclusiveScan(
- T input, ///< [in] Calling thread's input item.
- T &output, ///< [out] Calling thread's output item. May be aliased with \p input.
- ScanOp scan_op, ///< [in] Binary scan operator
- T &warp_aggregate) ///< [out] Warp-wide aggregate reduction of input items.
- * \brief Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. The \p warp_prefix_op value from thread-thread-lane<sub>0</sub> is applied to all scan outputs. Also computes the warp-wide \p warp_aggregate of all inputs for thread-thread-lane<sub>0</sub>.
- *
- * The \p warp_prefix_op functor must implement a member function <tt>T operator()(T warp_aggregate)</tt>.
- * The functor's input parameter \p warp_aggregate is the same value also returned by the scan operation.
- * The functor will be invoked by the entire warp of threads, however only the return value from
- * <em>lane</em><sub>0</sub> is applied as the threadblock-wide prefix. Can be stateful.
- *
- * Supports non-commutative scan operators.
- *
- * \smemreuse
- *
- * The code snippet below illustrates a single thread block of 32 threads (one warp) that progressively
- * computes an exclusive prefix max scan over multiple "tiles" of input using a
- * prefix functor to maintain a running total between block-wide scans. Each tile consists
- * of 32 integer items that are partitioned across the warp.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * // A stateful callback functor that maintains a running prefix to be applied
- * Suppose the input \p d_data is <tt>0, -1, 2, -3, 4, -5, ...</tt>.
- * The corresponding output for the first segment will be <tt>INT_MIN, 0, 0, 2, ..., 28, 30</tt>.
- * The output for the second segment will be <tt>30, 32, 32, 34, ..., 60, 62</tt>. Furthermore,
- * \p block_aggregate will be assigned \p 30 in all threads after the first scan, assigned \p 62 after the second
- * scan, etc.
- *
- * \tparam ScanOp <b>[inferred]</b> Binary scan operator type having member <tt>T operator()(const T &a, const T &b)</tt>
- * \tparam WarpPrefixOp <b>[inferred]</b> Call-back functor type having member <tt>T operator()(T warp_aggregate)</tt>
- */
- template <
- typename ScanOp,
- typename WarpPrefixOp>
- __device__ __forceinline__ void ExclusiveScan(
- T input, ///< [in] Calling thread's input item.
- T &output, ///< [out] Calling thread's output item. May be aliased with \p input.
- ScanOp scan_op, ///< [in] Binary scan operator
- T &warp_aggregate, ///< [out] <b>[<em>warp-lane</em><sub>0</sub> only]</b> Warp-wide aggregate reduction of input items (exclusive of the \p warp_prefix_op value).
- WarpPrefixOp &warp_prefix_op) ///< [in-out] <b>[<em>warp-lane</em><sub>0</sub> only]</b> Call-back functor for specifying a warp-wide prefix to be applied to all inputs.
, m_next_index(ViewAllocateWithoutInitializing("UnorderedMap next index"), capacity()+1) // +1 so that the *_at functions can always return a valid reference
+# error "It doesn't make sense to build this file unless the Kokkos::Serial device is enabled. If you see this message, it probably means that there is an error in Kokkos' CMake build infrastructure."
+# error "You must enable at least one of the following execution spaces in order to build this test: Kokkos::Threads, Kokkos::OpenMP, or Kokkos::Serial."
// The global parallel_reduce does not support vector_length other than 1 at the moment
if(policy.vector_length() > 1)
Impl::throw_runtime_exception( "Kokkos::parallel_reduce with a TeamPolicy using a vector length of greater than 1 is not currently supported for CUDA.");
// Functor's reduce memory, team scan memory, and team shared memory depend upon team size.
const int shmem_size_total = m_team_begin + m_shmem_begin + m_shmem_size ;
#if defined ( KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_CUDA )
- typedef Kokkos::Cuda DefaultExecutionSpace ;
+ typedef Cuda DefaultExecutionSpace ;
#elif defined ( KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_OPENMP )
typedef OpenMP DefaultExecutionSpace ;
#elif defined ( KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_THREADS )
typedef Threads DefaultExecutionSpace ;
#elif defined ( KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_SERIAL )
typedef Serial DefaultExecutionSpace ;
#else
# error "At least one of the following execution spaces must be defined in order to use Kokkos: Kokkos::Cuda, Kokkos::OpenMP, Kokkos::Serial, or Kokkos::Threads."
#endif
+#if defined ( KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_OPENMP )
+ typedef OpenMP DefaultHostExecutionSpace ;
+#elif defined ( KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_THREADS )
+ typedef Threads DefaultHostExecutionSpace ;
+#elif defined ( KOKKOS_HAVE_DEFAULT_DEVICE_TYPE_SERIAL )
+ typedef Serial DefaultHostExecutionSpace ;
+#elif defined ( KOKKOS_HAVE_OPENMP )
+ typedef OpenMP DefaultHostExecutionSpace ;
+#elif defined ( KOKKOS_HAVE_PTHREAD )
+ typedef Threads DefaultHostExecutionSpace ;
+#elif defined ( KOKKOS_HAVE_SERIAL )
+ typedef Serial DefaultHostExecutionSpace ;
+#else
+# error "At least one of the following execution spaces must be defined in order to use Kokkos: Kokkos::OpenMP, Kokkos::Serial, or Kokkos::Threads."
# error "At least one of the following host execution spaces must be defined: Kokkos::OpenMP, Kokkos::Serial, or Kokkos::Threads. You might be seeing this message if you disabled the Kokkos::Serial device explicitly using the Kokkos_ENABLE_Serial:BOOL=OFF CMake option, but did not enable any of the other host execution space devices."
if ((strncmp(arg[iarg],"--kokkos-threads",16) == 0) || (strncmp(arg[iarg],"--threads",9) == 0)) {
//Find the number of threads (expecting --threads=XX)
if (!((strncmp(arg[iarg],"--kokkos-threads=",17) == 0) || (strncmp(arg[iarg],"--threads=",10) == 0)))
Impl::throw_runtime_exception("Error: expecting an '=INT' after command line argument '--threads/--kokkos-threads'. Raised by Kokkos::initialize(int narg, char* argc[]).");
Impl::throw_runtime_exception("Error: expecting an '=INT' after command line argument '--threads/--kokkos-threads'. Raised by Kokkos::initialize(int narg, char* argc[]).");
if (!((strncmp(arg[iarg],"--kokkos-numa=",14) == 0) || (strncmp(arg[iarg],"--numa=",7) == 0)))
Impl::throw_runtime_exception("Error: expecting an '=INT' after command line argument '--numa/--kokkos-numa'. Raised by Kokkos::initialize(int narg, char* argc[]).");
Impl::throw_runtime_exception("Error: expecting an '=INT' after command line argument '--numa/--kokkos-numa'. Raised by Kokkos::initialize(int narg, char* argc[]).");
//Find the number of device (expecting --device=XX)
if (!((strncmp(arg[iarg],"--kokkos-device=",16) == 0) || (strncmp(arg[iarg],"--device=",9) == 0)))
Impl::throw_runtime_exception("Error: expecting an '=INT' after command line argument '--device/--kokkos-device'. Raised by Kokkos::initialize(int narg, char* argc[]).");
Impl::throw_runtime_exception("Error: expecting an '=INT' after command line argument '--device/--kokkos-device'. Raised by Kokkos::initialize(int narg, char* argc[]).");
//Find the number of device (expecting --device=XX)
if (!((strncmp(arg[iarg],"--kokkos-ndevices=",18) == 0) || (strncmp(arg[iarg],"--ndevices=",11) == 0)))
Impl::throw_runtime_exception("Error: expecting an '=INT[,INT]' after command line argument '--ndevices/--kokkos-ndevices'. Raised by Kokkos::initialize(int narg, char* argc[]).");
Impl::throw_runtime_exception("Error: expecting an integer number after command line argument '--kokkos-ndevices'. Raised by Kokkos::initialize(int narg, char* argc[]).");
Impl::throw_runtime_exception("Error: expecting an integer number after command line argument '--kokkos-ndevices=XX,'. Raised by Kokkos::initialize(int narg, char* argc[]).");
+\mainpage Trilinos/Kokkos: Shared-memory programming interface and computational kernels
+
+\section Kokkos_Intro Introduction
+
+The %Kokkos package has two main components. The first, sometimes
+called "%Kokkos Array" or just "%Kokkos," implements a
+performance-portable shared-memory parallel programming model and data
+containers. The second, called "%Kokkos Classic," consists of
+computational kernels that support the %Tpetra package.
+
+\section Kokkos_Kokkos The %Kokkos programming model
+
+%Kokkos implements a performance-portable shared-memory parallel
+programming model and data containers. It lets you write an algorithm
+once, and just change a template parameter to get the optimal data
+layout for your hardware. %Kokkos has back-ends for the following
+parallel programming models:
+
+- Kokkos::Threads: POSIX Threads (Pthreads)
+- Kokkos::OpenMP: OpenMP
+- Kokkos::Cuda: NVIDIA's CUDA programming model for graphics
+ processing units (GPUs)
+- Kokkos::Serial: No thread parallelism
+
+%Kokkos also has optimizations for shared-memory parallel systems with
+nonuniform memory access (NUMA). Its containers can hold data of any
+primitive ("plain old") data type (and some aggregate types). %Kokkos
+Array may be used as a stand-alone programming model.
+
+%Kokkos' parallel operations include the following:
+
+- parallel_for: a thread-parallel "for loop"
+- parallel_reduce: a thread-parallel reduction
+- parallel_scan: a thread-parallel prefix scan operation
+
+as well as expert-level platform-independent interfaces to thread
+"teams," per-team "shared memory," synchronization, and atomic update
+operations.
+
+%Kokkos' data containers include the following:
+
+- Kokkos::View: A multidimensional array suitable for thread-parallel
+ operations. Its layout (e.g., row-major or column-major) is
+ optimized by default for the particular thread-parallel device.
+- Kokkos::Vector: A drop-in replacement for std::vector that eases
+ porting from standard sequential C++ data structures to %Kokkos'
+ parallel data structures.
+- Kokkos::UnorderedMap: A parallel lookup table comparable in
+ functionality to std::unordered_map.
+
+%Kokkos also uses the above basic containers to implement higher-level
+data structures, like sparse graphs and matrices.
+
+A good place to start learning about %Kokkos would be <a href="http://trilinos.sandia.gov/events/trilinos_user_group_2013/presentations/2013-11-TUG-Kokkos-Tutorial.pdf">these tutorial slides</a> from the 2013 Trilinos Users' Group meeting.
+
+\section Kokkos_Classic %Kokkos Classic
+
+"%Kokkos Classic" consists of computational kernels that support the
+%Tpetra package. These kernels include sparse matrix-vector multiply,
+sparse triangular solve, Gauss-Seidel, and dense vector operations.
+They are templated on the type of objects (\c Scalar) on which they
+operate. This component was not meant to be visible to users; it is
+an implementation detail of the %Tpetra distributed linear algebra
+package.
+
+%Kokkos Classic also implements a shared-memory parallel programming
+model. This inspired and preceded the %Kokkos programming model
+described in the previous section. Users should consider the %Kokkos
+Classic programming model deprecated, and prefer the new %Kokkos